GPU-based RISC-V instruction set optimization method, system and device, and medium

By extending RISC-V instruction set and compiler optimization, the problem of RISC-V register competition in the GPU environment is solved, GPU performance and efficiency are improved, and parallel processing with lower power consumption is achieved.

CN119127308APending Publication Date: 2024-12-13SHANDONG INSPUR SCI RES INST CO LTD
View PDF 0 Cites 3 Cited by

Patent Information

Application Number
CN202411324804.7
Authority / Receiving Office
CN · China
Patent Type
Applications(China)
Current Assignee / Owner
Filing Date
2024-09-23
Publication Date
2024-12-13

AI Technical Summary

Technical Problem

The lack of optimization of the RISC-V instruction set in the GPU field has led to register competition when multiple threads frequently access general registers in the GPU environment, reducing execution efficiency and increasing power consumption.

Method used

Extend the RISC-V instruction set, add computational instructions that can access memory, optimize strategies through inline assembly and compiler optimization, reduce the use of general registers, optimize register allocation strategies, and support data operations directly from memory.

Benefits of technology

It improves the performance of RISC-V architecture GPUs when handling complex computing tasks, reduces register competition, reduces power consumption, improves parallel processing efficiency, and provides convenient optimization methods for advanced programming language developers.

✦ Generated by Eureka AI based on patent content.

Smart Images

  • Figure CN119127308A_ABST
    Figure CN119127308A_ABST
Patent Text Reader

Abstract

The invention relates to the technical field of processors of RISC-V architecture, and relates to a GPU-based RISC-V instruction set optimization method and device, and a medium. The method comprises the steps that an RISC-V instruction set is expanded, and a calculation instruction capable of accessing a memory is added; in an assembly programming stage, directly loading data from a memory in a calculation instruction; in a high-level programming language programming stage, through an inline assembly mode, an optimization instruction is packaged into an inline function, and a high-level language interface is provided for high-level language programming; and performing register allocation strategy optimization in the compiler, balancing the use frequency of the general register, and supporting and optimizing the inline function. According to the method, the performance of the RISC-V architecture GPU in processing complex computing tasks is remarkably improved. According to the optimization, the step of loading unnecessary data into the general register is reduced, so that the register competition is reduced, and the parallel processing efficiency is improved.
Need to check novelty before this filing date? Find Prior Art

Description

Technical Field

[0001] The present application relates to the technical field of RISC-V architecture processors, and more particularly to a GPU-based RISC-V instruction set optimization method, system, device, and medium. Background Art

[0002] Currently, RISC-V is an open source instruction set architecture developed on the basis of reduced instruction set (RISC), and its design concept aims to provide high flexibility and scalability to adapt to a variety of different applications and needs. As a modern, open source and fully open instruction set architecture, RISC-V has the characteristics of modular design, high-performance optimization, low-power design and good portability, which enables it to support a wide range of applications from simple microcontrollers to high-performance servers.

[0003] In recent years, with the intensification of global technology competition, especially the increasing demand for localization in the semiconductor field, the RISC-V instruction set has played an important role in the development of domestic GPUs. The open source characteristics and flexibility of the RISC-V architecture provide an ideal platform for domestic GPU design teams, greatly accelerating the development and industrialization of domestic GPUs.

[0004] However, although the RISC-V architecture has made significant progress in the CPU field, its application and optimization in the GPU architecture is relatively immature. Compared with traditional GPU architectures, such as NVIDIA's CUDA architecture and AMD's GCN architecture, RISC-V is still insufficiently optimized in the GPU field. These traditional GPU architectures have undergone years of optimization and accumulation, and have made a lot of special optimizations for graphics rendering and parallel computing, which makes them have significant advantages in performance and power efficiency.

[0005] In the GPU architecture, a notable feature is the execution mode of warps, where each warp usually contains multiple threads (such as 32 threads in NVIDIA CUDA). These threads need to share a set of general registers during execution. However, in the RISC-V instruction set, since only load and store instructions can access memory, other operation instructions cannot directly load or store data from memory. This leads to register competition in the GPU environment if multiple threads need to frequently access the same general register, which reduces the execution efficiency of the GPU and increases power consumption.

[0006] To address the above issues, traditional GPU architectures (such as CUDA) reduce the use of general registers by allowing computing instructions to load data directly from memory, thereby improving efficiency and reducing power consumption. However, the RISC-V instruction set does not currently support this optimization method, so it needs to be expanded and optimized to better meet the needs of the GPU architecture.

[0007] It should be noted that the information disclosed in the above background technology section is only used to enhance the understanding of the background of the present application, and therefore may include information that does not constitute the prior art known to ordinary technicians in the field. Summary of the invention

[0008] In order to provide a basic understanding of some aspects of the disclosed embodiments, a brief summary is given below. The summary is not an extensive review, nor is it intended to identify key / critical components or delineate the scope of protection of these embodiments, but rather serves as a prelude to the detailed description that follows.

[0009] The embodiments of the present disclosure provide a method and system, device, and medium for optimizing a RISC-V instruction set based on a GPU. By improving and optimizing the RISC-V instruction set, the RISC-V instruction set can better adapt to the characteristics of the GPU architecture, improve the performance of the GPU, and reduce power consumption, thereby promoting the further development of RISC-V in the GPU field.

[0010] In some embodiments, the method comprises:

[0011] Expand the RISC-V instruction set and add computing instructions that can access memory;

[0012] In the assembly programming stage, data is directly loaded from the memory in the calculation instructions to reduce the use of general registers;

[0013] In the high-level programming language programming stage, the optimization instructions are encapsulated into an inline function through inline assembly, and a high-level language interface is provided for high-level language programming; the register allocation strategy is optimized in the compiler to balance the usage frequency of general registers, support and optimize the inline function.

[0014] Preferably, the high-level programming language programming stage also includes increasing the number of general registers.

[0015] Preferably, the inline function in the high-level programming language programming stage is to encapsulate the optimization instruction into a function that can be called by the high-level programming language.

[0016] Preferably, the compiler optimization process includes:

[0017] Convert source code into an intermediate representation, which is a high-level abstraction used internally by the compiler to analyze and transform code;

[0018] Perform data flow analysis and control flow analysis at the intermediate representation level to track the definition and use of variables;

[0019] Build a use-definition chain to track the usage location of each variable and its corresponding definition location, and identify whether the load operation is only used in a single calculation instruction;

[0020] Refer to the instruction set manual of the target processor to identify the instructions that support direct memory operations and determine the memory operand requirements of these instructions;

[0021] Modify the instruction sequence in the intermediate representation, replacing separate load instructions with computation instructions that can operate directly from memory;

[0022] According to the optimized instruction sequence, the register allocation is updated to release the registers no longer needed due to the elimination of load operations for use by other variables or operations;

[0023] After the optimization is complete, a verification step is performed to ensure that the optimization operation has not changed the behavior of the original program.

[0024] Preferably, the verification step comprises running test cases or using formal verification techniques.

[0025] Preferably, tracking the definition and usage of the variable includes determining the definition timing, usage timing and life cycle of the variable.

[0026] In some embodiments, the system comprises:

[0027] An instruction execution unit, used to execute the GPU-based RISC-V instruction set optimization method to optimize the GPU-based RISC-V instruction set;

[0028] A compiler, integrated into the computer system or as an external module, for converting source code into an intermediate representation, performing data flow analysis and control flow analysis, and optimizing register allocation strategy;

[0029] An optimization engine, integrated in the compiler, is used to identify and call optimization instructions in the extended RISC-V instruction set, and to encapsulate the optimization instructions into inline functions.

[0030] Preferably, the compiler comprises:

[0031] A definition chain building block that tracks where each variable is used and where it is defined;

[0032] The instruction set manual checks the module to determine the instructions that support direct memory operations and the memory operand requirements of these instructions.

[0033] A register allocation update module, used for updating register allocation, and reallocating registers originally allocated to the eliminated load operation to other variables or operations;

[0034] The verification module is used to verify that the optimization has not changed the behavior of the program, including running test cases or using formal verification techniques.

[0035] In some embodiments, the device includes: a processor and a memory storing program instructions, and the processor is configured to execute the GPU-based RISC-V instruction set optimization method when running the program instructions.

[0036] The memory further includes: a storage area for storing an extended RISC-V instruction set; and a storage area for storing inline functions, wherein the inline functions encapsulate optimized instructions.

[0037] In some embodiments, the medium stores program instructions, and when the program instructions are run, they execute the GPU-based RISC-V instruction set optimization method.

[0038] The RISC-V instruction set optimization method, device, and medium based on GPU provided by the embodiments of the present disclosure can achieve the following technical effects:

[0039] The present invention significantly improves the performance of RISC-V architecture GPUs in processing complex computing tasks by optimizing the RISC-V instruction set based on GPUs. This optimization reduces the steps of loading unnecessary data into general registers, thereby reducing register competition and improving parallel processing efficiency.

[0040] The optimized instruction set reduces the reliance on general registers and avoids the extra power consumption caused by frequent register competition. This not only helps to extend the battery life of the device, but also conforms to the current technological trend of energy conservation and emission reduction.

[0041] Through inline assembly and optimization compiler strategy, the present invention provides a convenient optimization approach for developers using high-level programming languages, which not only reduces the difficulty of directly writing assembly code, but also improves the maintainability and portability of the code.

[0042] The present invention proposes a method for optimizing a compiler register allocation strategy, which further reduces register competition and improves resource utilization by balancing the usage frequency of general registers.

[0043] The extended RISC-V instruction set adds computing instructions that can access data directly from memory. This design reduces the number of times data is transferred between memory and registers, thereby improving data processing efficiency.

[0044] The above general description and the following description are exemplary and explanatory only and are not intended to limit the present application. BRIEF DESCRIPTION OF THE DRAWINGS

[0045] One or more embodiments are exemplarily described by corresponding drawings, which do not limit the embodiments. Elements with the same reference numerals in the drawings are shown as similar elements, and the drawings do not constitute a scale limitation, and wherein:

[0046] Figure 1 It is a schematic flow chart of the method of the present invention;

[0047] Figure 2 It is a flowchart of a compiler calling an optimization instruction provided by an embodiment of the present disclosure;

[0048] Figure 3 It is a schematic diagram of an optimization device for the GPU-based RISC-V instruction set provided in an embodiment of the present disclosure. DETAILED DESCRIPTION

[0049] In order to be able to understand the features and technical contents of the embodiments of the present disclosure in more detail, the implementation of the embodiments of the present disclosure is described in detail below in conjunction with the accompanying drawings. The attached drawings are for reference only and are not used to limit the embodiments of the present disclosure. In the following technical description, for the convenience of explanation, a full understanding of the disclosed embodiments is provided through multiple details. However, one or more embodiments can still be implemented without these details. In other cases, to simplify the drawings, well-known structures and devices can be simplified for display.

[0050] The terms "first", "second", etc. in the specification and claims of the embodiments of the present disclosure and the above-mentioned drawings are used to distinguish similar objects, and are not necessarily used to describe a specific order or sequence. It should be understood that the terms used in this way can be interchanged where appropriate, so as to describe the embodiments of the embodiments of the present disclosure described herein. In addition, the terms "including" and "having" and any variations thereof are intended to cover non-exclusive inclusions.

[0051] Example 1

[0052] Combination Figure 1 , a GPU-based RISC-V instruction set optimization method, the method comprising:

[0053] S1: Expand the RISC-V instruction set and add computing instructions that can access memory.

[0054] S2: In the assembly programming stage, data is directly loaded from the memory in the calculation instructions to reduce the use of general registers.

[0055] S3: In the high-level programming language programming stage, the optimization instructions are encapsulated into an inline function through inline assembly, and a high-level language interface is provided for high-level language programming; register allocation strategy optimization is performed in the compiler to balance the usage frequency of general registers, and the inline function is supported and optimized.

[0056] As a refinement of the above embodiment, the high-level programming language programming stage also includes increasing the number of general registers.

[0057] As a refinement of the above embodiment, the inline function in the high-level programming language programming stage is to encapsulate the optimization instruction into a function that can be called by the high-level programming language.

[0058] As a refinement of the above embodiment, Figure 2 As shown, the compiler optimization process includes:

[0059] S301: The compiler first converts the source code into an intermediate representation, which is a high-level abstract form used by the compiler to analyze and transform the code. At the intermediate representation level, the compiler performs data flow analysis and control flow analysis on the code.

[0060] S302: The compiler performs data flow analysis to track the definition and use of variables, including determining when variables are defined, when they are used, and their life cycles.

[0061] S303: The compiler constructs a definition chain to track the location where each variable is used and its definition location. In this way, the compiler can determine whether the result of a load operation is only used in a single calculation instruction.

[0062] S304: The compiler checks the instruction set manual of the target processor to determine which instructions support direct memory operations and the requirements of these instructions on memory operands (such as alignment, access size, etc.).

[0063] S305: The compiler modifies the original instruction sequence in the intermediate representation, replacing the separate load instructions with calculation instructions that can be directly operated from the memory.

[0064] S306: If the load operation is eliminated, the register originally allocated to this operation may no longer be needed. The compiler will update the register allocation and reallocate the register to other variables or operations.

[0065] S307: After optimization, the compiler usually performs verification to ensure that the optimization has not changed the behavior of the program. This may include running test cases or using formal verification techniques.

[0066] As a refinement of the above embodiment, the verification step includes running test cases or using formal verification technology.

[0067] As a refinement of the above embodiment, tracking the definition and usage of variables includes determining the definition timing, usage timing and life cycle of the variables.

[0068] The present invention is described below in conjunction with specific RISC-V architecture code:

[0069] In the RISC-V architecture, write the following code:

[0070] lw a1,rd; / / rd is a 32-bit address. This address in the memory stores two single-precision floating-point values ​​required for calculation. Load this address into general register a1

[0071] flw f1,0(a1) / / Load the floating point value into register f1 from the address pointed to by a1

[0072] flw f2,4(a1) / / Load another floating point value into register f2 by shifting 4 from the address pointed to by a1

[0073] fmul.s f3,f1,f2 / / Multiply the floating point numbers in f1 and f2 and save the result to register f3

[0074] fm.s f3,f1,f3 / / Merge the floating point numbers in f1 and f3 and save the result in f3

[0075] In the above code, the floating point number is loaded from the memory to the general register f2. In the subsequent calculation process, the general register f2 is only called once. The above situation is very common in RISC-V CPU programming, because the RISC-V instruction set manual requires that only load and store instructions can access the memory, that is, it is not allowed to directly load or store data from the memory in the operation instruction. Even if the subsequent calculation process is only used once, the load instruction must be used first to load the data from the memory to the general register.

[0076] The above situation will not cause problems in a single-threaded or small-threaded CPU, but in a GPU, a thread warp usually contains 32 threads, and the 32 threads share a set of general registers. If these 32 threads need to access the same general register, the scheduler needs to ensure that these threads access this general register one by one in a certain order. This general register competition situation greatly reduces the efficiency of the GPU, uses registers unnecessarily, and increases the power consumption of the GPU.

[0077] For the above situation, if we can avoid loading the floating-point number into the general register f2 and instead load it directly from the memory in the subsequent calculation instructions, we can reduce the use of general registers, thereby reducing usage conflicts, improving the efficiency of the GPU and reducing the power consumption of the GPU.

[0078] In the traditional GPU architecture, taking NVIDIA's CUDA architecture as an example, the same function code above is written as follows:

[0079] ld.global.f32%f1,rd / / Load a floating point number from address rd of global memory to general register %f1

[0080] ld.global.f32%f2,(rd+4) / / Load a floating point number from address rd+4 of global memory to general register %f2

[0081] mul.ftz.f32 %f3, %f1, %f2 / / The value in general register %f1 is multiplied by the value in %f2 and the result is saved in general register %f3

[0082] fma.rn.ftz.f32 %f3, %f1, %f3, %f3 / / fused multiplication and addition, the result is saved to %f3

[0083] The above code has the same problem as the RISC-V code. The floating point number is loaded from the memory to the general register f2. In the subsequent calculation process, the general register f2 is called once. Then in the CUDA architecture, we can use the following code to optimize the above code.

[0084] ld.global.f32%f1,rd / / Load a floating point number from address rd of global memory to general register %f1

[0085] mul.ftz.f32%f3,rd+4,%f1 / / Load a floating point number from address rd+4 of global memory and multiply it by the value in %f1, and save the result to %f3

[0086] fma.rn.ftz.f32 %f3, %f1, %f3, %f3 / / fused multiplication and addition, the result is saved to %f3

[0087] In the above code, the floating point number is not loaded into the general register f2 in advance, but is directly loaded from the memory in the multiplication instruction. This optimization greatly reduces the use of general registers.

[0088] To this end, we should expand the RISC-V instruction set accordingly and add computing instructions that can access memory, including but not limited to the following instructions, where rd represents the memory address:

[0089] fadd.s f2,f1,rd / / Load a floating point number from memory address rd and add it to the floating point number in f1

[0090] fsub.s f2,f1,rd / / Load a floating point number from memory address rd and subtract it from the floating point number in f1

[0091] fmul.s f2,f1,rd / / Load a floating point number from memory address rd and multiply it by the floating point number in f1

[0092] fdiv.s f2,f1,rd / / Load a floating point number from memory address rd and divide it by the floating point number in f1

[0093] It should be noted that the technical solution of this embodiment significantly improves the performance of the RISC-V architecture GPU when processing complex computing tasks. This optimization reduces the steps of loading unnecessary data into general registers, thereby reducing register competition and improving parallel processing efficiency. The optimized instruction set reduces the dependence on general registers and avoids the additional power consumption caused by frequent register competition. Through inline assembly and optimization compiler strategies, the present invention provides a convenient optimization path for developers using high-level programming languages. This not only reduces the difficulty of directly writing assembly code, but also improves the maintainability and portability of the code.

[0094] Example 2

[0095] In some embodiments, the system comprises:

[0096] An instruction execution unit, used to execute the GPU-based RISC-V instruction set optimization method to optimize the GPU-based RISC-V instruction set;

[0097] A compiler, integrated into the computer system or as an external module, for converting source code into an intermediate representation, performing data flow analysis and control flow analysis, and optimizing register allocation strategy;

[0098] An optimization engine, integrated in the compiler, is used to identify and call optimization instructions in the extended RISC-V instruction set, and to encapsulate the optimization instructions into inline functions.

[0099] As a refinement of the above embodiment, the compiler includes:

[0100] A definition chain building block that tracks where each variable is used and where it is defined;

[0101] The instruction set manual checks the module to determine the instructions that support direct memory operations and the memory operand requirements of these instructions.

[0102] A register allocation update module, used for updating register allocation, and reallocating registers originally allocated to the eliminated load operation to other variables or operations;

[0103] The verification module is used to verify that the optimization has not changed the behavior of the program, including running test cases or using formal verification techniques.

[0104] In some embodiments, the device includes: a processor and a memory storing program instructions, and the processor is configured to execute the GPU-based RISC-V instruction set optimization method when running the program instructions.

[0105] Combination Figure 3 As shown, the embodiment of the present disclosure provides an optimization device 300 based on the RISC-V instruction set of a GPU, including a processor (processor) 304 and a memory (memory) 301. Optionally, the device may also include a communication interface (Communication Interface) 302 and a bus 303. Among them, the processor 304, the communication interface 302, and the memory 301 can communicate with each other through the bus 303. The communication interface 302 can be used for information transmission. The processor 304 can call the logic instructions in the memory 301 to execute the optimization method of the RISC-V instruction set based on the GPU of the above embodiment.

[0106] In addition, the logic instructions in the memory 301 described above can be implemented in the form of software functional units and can be stored in a computer-readable storage medium when sold or used as an independent product.

[0107] The memory 301 is a computer-readable storage medium that can be used to store software programs and computer executable programs, such as program instructions / modules corresponding to the method in the embodiment of the present disclosure. The processor 304 executes the functional application and data processing by running the program instructions / modules stored in the memory 301, that is, the optimization method of the RISC-V instruction set based on the GPU in the above embodiment is implemented.

[0108] The memory 301 may include a program storage area and a data storage area, wherein the program storage area may store an operating system and an application required for at least one function; the data storage area may store data created according to the use of the terminal device, etc. In addition, the memory 301 may include a high-speed random access memory and may also include a non-volatile memory.

[0109] An embodiment of the present disclosure provides a computer-readable storage medium storing computer-executable instructions, wherein the computer-executable instructions are configured to execute the above-mentioned GPU-based RISC-V instruction set optimization method.

[0110] The computer-readable storage medium mentioned above may be a transient computer-readable storage medium or a non-transitory computer-readable storage medium.

[0111] The technical solution of the embodiment of the present disclosure can be embodied in the form of a software product, which is stored in a storage medium and includes one or more instructions for causing a computer device (which may be a personal computer, a server, or a network device, etc.) to execute all or part of the steps of the method described in the embodiment of the present disclosure. The aforementioned storage medium may be a non-transient storage medium, including: a USB flash drive, a mobile hard disk, a read-only memory (ROM), a random access memory (RAM), a disk or an optical disk, and other media that can store program codes, or a transient storage medium.

[0112] The above description and the accompanying drawings fully illustrate the embodiments of the present disclosure so that those skilled in the art can practice them. Other embodiments may include structural, logical, electrical, process and other changes. The embodiments represent only possible changes. Unless explicitly required, separate components and functions are optional, and the order of operation may vary. The parts and features of some embodiments may be included in or replace the parts and features of other embodiments. Moreover, the words used in this application are only used to describe the embodiments and are not used to limit the claims. As used in the description of the embodiments and the claims, unless the context clearly indicates, the singular forms of "a", "an" and "the" are intended to include plural forms as well. Similarly, the term "and / or" as used in this application refers to any and all possible combinations of listings containing one or more associated ones. In addition, when used in the present application, the term "comprise" and its variants "comprises" and / or comprising refer to the presence of stated features, wholes, steps, operations, elements, and / or components, but do not exclude the presence or addition of one or more other features, wholes, steps, operations, elements, components and / or groups thereof. In the absence of further restrictions, the elements defined by the sentence "comprising a ..." do not exclude the presence of other identical elements in the process, method or device comprising the elements. In this article, each embodiment may focus on the differences from other embodiments, and the same and similar parts between the embodiments may refer to each other. For the methods, products, etc. disclosed in the embodiments, if they correspond to the method part disclosed in the embodiments, then the relevant parts can refer to the description of the method part.

[0113] Those skilled in the art will appreciate that the units and algorithm steps of each example described in conjunction with the embodiments disclosed herein can be implemented in electronic hardware, or a combination of computer software and electronic hardware. Whether these functions are performed in hardware or software may depend on the specific application and design constraints of the technical solution. The technicians may use different methods to implement the described functions for each specific application, but such implementation should not be considered to be beyond the scope of the embodiments of the present disclosure. The technicians may clearly understand that for the convenience and simplicity of description, the specific working processes of the systems, devices and units described above can refer to the corresponding processes in the aforementioned method embodiments, and will not be repeated here.

[0114] In the embodiments disclosed herein, the disclosed methods and products (including but not limited to devices, equipment, etc.) can be implemented in other ways. For example, the device embodiments described above are only schematic. For example, the division of the units can be only a logical function division. There may be other division methods in actual implementation, such as multiple units or components can be combined or integrated into another system, or some features can be ignored or not executed. In addition, the coupling or direct coupling or communication connection between each other shown or discussed can be through some interfaces, indirect coupling or communication connection of devices or units, which can be electrical, mechanical or other forms. The units described as separate components may or may not be physically separated, and the components displayed as units may or may not be physical units, that is, they may be located in one place, or they may be distributed on multiple network units. Some or all of the units may be selected according to actual needs to implement this embodiment. In addition, each functional unit in the embodiment of the present disclosure may be integrated in a processing unit, or each unit may exist physically alone, or two or more units may be integrated in one unit.

[0115] The flowchart and block diagram in the accompanying drawings show the possible architecture, function and operation of the system, method and computer program product according to the embodiment of the present disclosure. In this regard, each box in the flowchart or block diagram can represent a module, a program segment or a part of the code, and the module, the program segment or a part of the code contains one or more executable instructions for realizing the specified logical function. In some alternative implementations, the functions marked in the box can also occur in a different order from the order marked in the accompanying drawings. For example, two consecutive boxes can actually be executed substantially in parallel, and they can sometimes be executed in the opposite order, which can depend on the functions involved. In the description corresponding to the flowchart and the block diagram in the accompanying drawings, the operations or steps corresponding to different boxes can also occur in a different order from the order disclosed in the description, and sometimes there is no specific order between different operations or steps. For example, two consecutive operations or steps can actually be executed substantially in parallel, and they can sometimes be executed in the opposite order, which can depend on the functions involved. Each block in the block diagrams and / or flowcharts, and combinations of blocks in the block diagrams and / or flowcharts, may be implemented by a dedicated hardware-based system that performs the specified functions or actions, or may be implemented by a combination of dedicated hardware and computer instructions.

Claims

1. A GPU-based RISC-V instruction set optimization method, characterized in that: The method comprises the following steps: Expand the RISC-V instruction set and add computing instructions that can access memory; In the assembly programming stage, data is directly loaded from the memory in the calculation instructions to reduce the use of general registers; In the high-level programming language programming stage, the optimization instructions are encapsulated into an inline function through inline assembly, and a high-level language interface is provided for high-level language programming; the register allocation strategy is optimized in the compiler to balance the usage frequency of general registers, support and optimize the inline function.

2. The GPU-based RISC-V instruction set optimization method according to claim 1, characterized in that: The high-level programming language programming stage also includes increasing the number of general registers.

3. The GPU-based RISC-V instruction set optimization method according to claim 1, characterized in that: The inline function described in the high-level programming language programming stage is to encapsulate the optimization instruction into a function that can be called by the high-level programming language.

4. The GPU-based RISC-V instruction set optimization method according to claim 1, characterized in that: The compiler optimization process includes: Convert source code into an intermediate representation, which is a high-level abstraction used internally by the compiler to analyze and transform code; Perform data flow analysis and control flow analysis at the intermediate representation level to track the definition and use of variables; Build a use-definition chain to track the usage location of each variable and its corresponding definition location, and identify whether the load operation is only used in a single calculation instruction; Refer to the instruction set manual of the target processor to identify the instructions that support direct memory operations and determine the memory operand requirements of these instructions; Modify the instruction sequence in the intermediate representation, replacing separate load instructions with computation instructions that can operate directly from memory; According to the optimized instruction sequence, the register allocation is updated to release the registers no longer needed due to the elimination of load operations for use by other variables or operations; After the optimization is complete, a verification step is performed to ensure that the optimization operation has not changed the behavior of the original program.

5. The GPU-based RISC-V instruction set optimization method according to claim 4, characterized in that: The verification step includes running test cases or using formal verification techniques.

6. The GPU-based RISC-V instruction set optimization method according to claim 1, characterized in that: Tracking the definition and use of variables includes determining when the variable is defined, when it is used, and its life cycle.

7. A GPU-based RISC-V instruction set optimization system, characterized in that: The system includes: An instruction execution unit, used to execute the GPU-based RISC-V instruction set optimization method to optimize the GPU-based RISC-V instruction set; A compiler, integrated into the computer system or as an external module, for converting source code into an intermediate representation, performing data flow analysis and control flow analysis, and optimizing register allocation strategy; An optimization engine, integrated in the compiler, is used to identify and call optimization instructions in the extended RISC-V instruction set, and to encapsulate the optimization instructions into inline functions.

8. The GPU-based RISC-V instruction set optimization system according to claim 7, characterized in that: The compiler comprises: A definition chain building block that tracks where each variable is used and where it is defined; An instruction set manual inspection module to determine the instructions that support direct memory operations and the memory operand requirements of these instructions; A register allocation update module, used for updating register allocation, and reallocating registers originally allocated to the eliminated load operation to other variables or operations; The verification module is used to verify that the optimization has not changed the behavior of the program, including running test cases or using formal verification techniques.

9. A GPU-based RISC-V instruction set optimization device, comprising a processor and a memory storing program instructions, characterized in that: The processor is configured to execute the GPU-based RISC-V instruction set optimization method as described in any one of claims 1 to 8 when running the program instructions.

10. A storage medium storing program instructions, characterized in that: When the program instructions are run, they execute the GPU-based RISC-V instruction set optimization method as described in any one of claims 1 to 8.

Citation Information

Cited By

  • RISC-V instruction acceleration method, system and device and storage medium

    CN119415155A

  • SIMD instruction arrangement method based on RISC-V architecture

    CN120447973A

  • Memory access optimization method and device for ensemble communication operator and computing equipment

    CN121210153A