An instruction set translation method for GPU heterogeneous environments based on simulation software

Abstract

This invention relates to the field of instruction set conversion technology and discloses an instruction set conversion method for GPU heterogeneous environments based on simulation software. The invention accurately captures key instruction flows in simulation tasks through dynamic instrumentation, providing a solid foundation for subsequent processing. Through a classification mechanism, the original instruction flow is divided into multiple instruction blocks to be converted based on instruction semantic features and hardware adaptability, creating conditions for parallel conversion. A lightweight front-end translator efficiently converts the classified instruction blocks into architecture-independent intermediate representations, preserving the parallelism and data dependencies between instructions. Deep optimization is performed on the intermediate representation instructions, including key steps such as instruction selection, register allocation, instruction reassembly, SIMT mapping, and instruction encoding. The target GPU thread blocks accelerate the translation and generation of native instructions for the target GPU architecture. This invention not only shortens the instruction conversion cycle but also improves the accuracy of the conversion results and execution efficiency.

Description

Technical Field

[0001] This invention relates to the field of instruction set conversion technology, and discloses an instruction set conversion method for a GPU heterogeneous environment based on simulation software. Background Technology

[0002] Current heterogeneous computing systems (such as CPU-GPU architectures) face the challenge of instruction set incompatibility when running legacy programs or cross-platform applications. Traditional solutions rely on static binary translation or hardware customization optimization, which has the following problems: First, low translation efficiency, as static translation schemes cannot utilize the parallel computing capabilities of GPUs to accelerate the translation process, resulting in limited translation throughput; second, poor dynamic adaptability, lacking runtime feedback mechanisms in scenarios with varying hardware configurations and dynamic load changes, making it impossible to dynamically adjust translation strategies based on actual execution conditions; third, high resource overhead, as frequent data copying between the CPU and GPU occurs during translation, and in multiphysics simulations (such as multi-stage compressor LES scenarios), existing QEMUs often have CPU utilization exceeding 90% and GPU idle rate exceeding 60% during PTX->GCN conversion, severely wasting hardware resources; fourth, insufficient scalability, making it difficult to adapt to the diverse instruction sets of emerging GPU architectures (such as NVIDIA CUDA PTX, AMD GCN, and domestic GPU instruction sets). Summary of the Invention

[0003] The purpose of this invention is to provide an instruction set conversion method for GPU heterogeneous environments based on simulation software, which can improve translation efficiency and reduce translation latency.

[0004] To achieve the above-mentioned technical effects, the technical solution adopted by the present invention is as follows:

[0005] A method for instruction set translation in a GPU heterogeneous environment based on simulation software, comprising:

[0006] The simulation software uses a dynamic instrumentation method to intercept the original instruction stream submitted by the simulation task program to the original GPU / CPU for execution and marks the instruction blocks to be converted; then, based on the instruction semantic features and hardware resource matching degree, the instruction blocks to be converted are classified and their priorities are determined.

[0007] A lightweight front-end translator is used to convert the instruction block to be translated into an architecture-independent intermediate representation instruction, while preserving the SIMT parallelism and data dependencies of the instruction block to be translated. Then, a translation task is defined, which includes using the target back-end optimizer to sequentially complete instruction selection, register allocation, instruction reassembly, SIMT mapping and instruction encoding on the intermediate representation instruction to obtain the native instructions of the target GPU architecture.

[0008] According to the priority order from high to low, the target GPU thread block is allocated to the intermediate representation instructions corresponding to the instruction block to be converted according to the preset scheduling strategy. Then, the translation tasks of the batch intermediate representation instructions are mapped to the corresponding target GPU thread blocks, and batch processing is carried out using SIMT parallelism to translate and obtain the native instructions of the target GPU architecture, and store them in the GPU video memory output buffer.

[0009] The target GPU executes the translated native instructions, collects performance metrics during the execution of the target GPU, and dynamically adjusts the subsequent translation strategy according to a preset feedback mechanism;

[0010] The translated native instruction sequence is stored according to a multi-level translation caching strategy, and adaptive cache management is performed according to the incremental cache update and eviction strategy.

[0011] Furthermore, the method for classifying the instruction blocks to be converted based on instruction semantic features and hardware resource matching degree is as follows:

[0012] Calculate the ratio of computational instructions to memory access instructions in the instruction block to be converted;

[0013] If the proportion of computational instructions in the instruction block to be converted exceeds a first preset threshold, the instruction block to be converted will be classified as a computationally intensive instruction block.

[0014] If the proportion of memory access instructions in the instruction block to be converted exceeds the second preset threshold, the instruction block to be converted will be classified as a memory access intensive instruction block.

[0015] If the proportion of computational instructions in the instruction block to be converted does not exceed the first preset threshold and the proportion of memory access instructions does not exceed the second preset threshold, then the instruction block to be converted is classified as a mixed instruction block.

[0016] Furthermore, methods for using a lightweight front-end translator to convert instruction blocks into architecture-independent intermediate representation instructions while preserving the SIMT parallelism and data dependencies of the instruction blocks include:

[0017] Each block of instructions to be converted is parsed into a structured PTX instruction sequence using a lexical analyzer;

[0018] The semantic analyzer is used to convert each PTX instruction into a corresponding intermediate representation instruction according to a predefined mapping rule, resulting in an intermediate representation instruction sequence for each instruction block to be converted.

[0019] Analyze and extract the SIMT parallel structure information of the instruction block to be converted, and append the obtained SIMT parallel structure information to the corresponding intermediate representation instruction sequence in the form of metadata;

[0020] Traverse the intermediate representation instruction sequence corresponding to the instruction block to be converted, identify the data dependency relationship between each intermediate representation instruction and other intermediate representation instructions in the sequence, and construct a data flow graph between instructions.

[0021] Furthermore, the data dependencies include RAW dependencies (read-after-write dependencies), WAR dependencies (read-after-write dependencies), WAW dependencies (write-after-write dependencies), and memory dependencies.

[0022] Furthermore, the method for obtaining the native instructions of the target GPU architecture by using the target back-end optimizer to sequentially perform instruction selection, register allocation, instruction reassembly, SIMT mapping, and instruction encoding on the intermediate representation instructions includes:

[0023] According to the instruction set specification of the target GPU architecture, each intermediate representation instruction is mapped to the corresponding target GPU architecture instruction to obtain the target GPU architecture instruction sequence;

[0024] A graph coloring algorithm is used to allocate physical registers for the virtual registers in the intermediate representation instructions;

[0025] Based on the scheduling characteristics of the parallel execution thread set in the target GPU architecture, the instruction sequence of the target GPU architecture is rearranged, and then SIMT mapping is performed. The SIMT-mapped instruction sequence of the target GPU architecture is packaged into directly executable machine code according to the binary encoding format of the target GPU, thus obtaining the native instruction sequence that can be executed by the target AMD GPU.

[0026] Furthermore, when using SIMT parallelism for batch processing, rapid synchronization of intermediate translation states is achieved between GPU thread blocks through shared memory and atomic operations, including the following steps:

[0027] A dedicated area is allocated in the shared memory of the target GPU to store the intermediate translation state;

[0028] For the first data that only needs to be read, each GPU thread block reads it in parallel from the GPU shared memory without synchronization between GPU thread blocks; for the second data that needs to be updated by multiple GPU thread blocks, atomic operations are used for updating.

[0029] At the switching points between translation stages, the intermediate translation state is synchronized between GPU thread blocks using the thread block synchronization barrier instruction.

[0030] Furthermore, the intermediate translation state refers to the temporary data and status information generated at each translation stage during the parallel translation process involving multiple GPU thread blocks.

[0031] Furthermore, the multi-level translation caching strategy includes an L1 cache and an L2 cache; the L1 cache is an instruction fragment level cache, used to store the translation results of frequently occurring basic blocks, using the hash value of the instruction block content to be converted as the index key and the corresponding target GPU architecture binary code as the value; the L2 cache is a function level cache, used to store the translation results of complete kernel functions, using the kernel function signature as the index key.

[0032] Compared with the prior art, the beneficial effects of this invention are:

[0033] This invention accurately captures key instruction flows in simulation tasks using a dynamic instrumentation method, providing a solid foundation for subsequent processing. Then, through a classification mechanism, the original instruction flow is divided into multiple instruction blocks to be converted based on instruction semantic features and hardware adaptability, creating conditions for parallel conversion. Next, a lightweight front-end translator efficiently converts the classified instruction blocks into architecture-independent intermediate representations, preserving parallelism and data dependencies between instructions while providing flexible space for subsequent optimization. Finally, by defining detailed translation tasks and leveraging a target back-end optimizer to deeply optimize the intermediate representation instructions, including key steps such as instruction selection, register allocation, instruction reassembly, SIMT mapping, and instruction encoding, the target GPU thread blocks accelerate the translation to generate native instructions for the target GPU architecture. This invention not only shortens the instruction conversion cycle but also improves the accuracy and execution efficiency of the conversion results. Attached Figure Description

[0034] Figure 1 This is a flowchart of the instruction set conversion method for a GPU heterogeneous environment based on simulation software in the embodiment. Detailed Implementation

[0035] The present invention will now be described in further detail with reference to the embodiments and accompanying drawings. However, this should not be construed as limiting the scope of the above-described subject matter of the present invention to the following embodiments; all technologies implemented based on the content of the present invention fall within the scope of the present invention.

[0036] See Figure 1 This invention provides an instruction set conversion method for GPU heterogeneous environments based on simulation software, comprising:

[0037] The simulation software uses a dynamic instrumentation method to intercept the original instruction stream submitted by the simulation task program to the original GPU / CPU for execution and marks the instruction blocks to be converted; then, based on the instruction semantic features and hardware resource matching degree, the instruction blocks to be converted are classified and their priorities are determined.

[0038] A lightweight front-end translator is used to convert the instruction block to be converted into intermediate representation instructions, while preserving the SIMT parallelism and data dependencies of the instruction block to be converted. Then, a translation task is defined, which includes using the target back-end optimizer to sequentially complete instruction selection, register allocation, instruction reassembly, SIMT mapping and instruction encoding of the intermediate representation instructions to obtain the native instructions of the target GPU architecture.

[0039] According to the priority order from high to low, target GPU thread blocks are allocated to the intermediate representation instructions corresponding to the instruction block to be converted according to the preset scheduling strategy. Then, the translation tasks of batch intermediate representation instructions are mapped to the corresponding target GPU thread blocks, and batch processing is carried out using SIMT parallelism to translate and obtain the native instructions of the target GPU architecture.

[0040] This invention accurately captures key instruction flows in simulation tasks using a dynamic instrumentation method, providing a solid foundation for subsequent processing. Then, through a classification mechanism, the original instruction flow is divided into multiple instruction blocks to be converted based on instruction semantic features and hardware adaptability, creating conditions for parallel conversion. Next, a lightweight front-end translator efficiently converts the classified instruction blocks into architecture-independent intermediate representations, preserving parallelism and data dependencies between instructions while providing flexible space for subsequent optimization. Finally, by defining detailed translation tasks and leveraging a target back-end optimizer to deeply optimize the intermediate representation instructions, including key steps such as instruction selection, register allocation, instruction reassembly, SIMT mapping, and instruction encoding, the target GPU thread blocks accelerate the translation to generate native instructions for the target GPU architecture. This invention not only shortens the instruction conversion cycle but also improves the accuracy and execution efficiency of the conversion results.

[0041] Example

[0042] This embodiment uses the conversion from CUDA to AMD GCN instruction set as an example to illustrate the implementation steps of this method:

[0043] Step 1: Instruction capture and classification: The simulation software uses dynamic instrumentation to intercept the original instruction stream submitted by the simulation task program to the original GPU / CPU for execution, and marks the instruction blocks to be converted; then, based on the instruction semantics and hardware resource matching degree, the instruction blocks to be converted are classified and their priorities are determined.

[0044] Step 1.1: Dynamic instrumentation is a method for intercepting and modifying the instruction stream during application runtime. This invention utilizes dynamic instrumentation to capture and mark GPU-related instruction blocks to be transformed. The specific implementation process is as follows:

[0045] First, an API hook is set in the GPU driver layer or runtime library layer of the CUDA application. When the CUDA application calls the relevant functions of the GPU runtime library, the API hook intercepts the call and redirects the control flow to a translation framework based on the instruction set conversion method of this invention. The translation framework then parses the parameter information of the intercepted API and extracts the GPU instruction stream to be executed, i.e., the instruction block to be converted, which is the PTX instruction stream for the CUDA program. The parameter information of the intercepted API includes kernel function pointers, startup configuration parameters (grid dimension, thread block dimension), kernel parameter lists, etc. The instruction block to be converted includes kernel startup instructions, memory operation instructions, PTX calculation instructions, PTX memory access instructions, and synchronization instructions.

[0046] The kernel startup instructions, such as cudaLaunchKernel in CUDA and clEnqueueNDRangeKernel in OpenCL, are responsible for submitting the computation kernel to the GPU for execution. The memory operation instructions include cudaMalloc (memory allocation), cudaMemcpy (data transfer), and cudaFree (memory release), which manage GPU memory resources. The PTX computation instructions include arithmetic logic instructions, control flow instructions, and special function instructions. Arithmetic logic instructions include add, mul, mad, and fma; control flow instructions include bra, call, and ret; and special function instructions include sin, cos, and rsqrt. These instructions constitute the core of the GPU computation kernel. The PTX memory access instructions include ld.global and st.global (global access), ld.shared and st.shared (shared memory access), and tex (texture sampling), which handle data exchange between the GPU kernel and various levels of memory. The synchronization instructions include bar.sync (thread synchronization) and membar (membar), which coordinate the execution order and memory consistency between GPU thread blocks.

[0047] Finally, the extracted instruction blocks to be converted are marked with key attributes such as instruction type, operands, target register, and memory address space, so that the translation engine can perform targeted processing in the future.

[0048] Step 1.2: Classify and prioritize the instruction blocks to be converted based on instruction semantic features and hardware resource matching degree.

[0049] First, the instruction blocks to be converted are classified based on instruction semantic features and hardware resource matching degree. In terms of instruction semantic features, the distribution of various opcodes in the instruction blocks to be converted is analyzed by static analysis, and the proportion of computational instructions (i.e., ALU arithmetic logic operations) and memory access instructions (i.e., load / store memory read and write operations) in the instruction blocks to be converted is statistically analyzed. In terms of hardware resource matching degree, the degree of adaptation between various instruction blocks to be converted and available hardware resources is evaluated based on the hardware configuration parameters of the target GPU, including the number of computing units (SM / CU number), memory bandwidth, cache capacity at each level, register file size, etc.: (1) If the proportion of computational instructions in the instruction block to be converted exceeds the first preset threshold, such as 70%, the instruction block to be converted is classified as a computationally intensive instruction block, and it is preferentially allocated to GPUs with sufficient computing units (i.e. GPUs with a large number of SMs or CUs) during translation. During the translation process, the focus is on register allocation optimization to maximize computational throughput. (2) If the proportion of memory access instructions in the instruction block to be translated exceeds the second preset threshold, such as 50%, then the instruction block to be translated is classified as a memory-intensive instruction block. During translation, it is preferentially allocated to GPUs with higher video memory bandwidth. The translation process focuses on memory coalescing and shared memory bank conflict elimination to make full use of video memory bandwidth. Through the classification strategy of hardware and software collaboration between instruction semantic features and hardware resource matching, the translated instruction block to be translated can be executed on suitable GPU hardware, thereby maximizing execution efficiency.

[0050] (3) If the proportion of computation-type instructions in the instruction block to be converted does not exceed the first preset threshold and the proportion of memory access instructions does not exceed the second preset threshold, then the instruction block to be converted is classified as a hybrid instruction block, and a general translation strategy is used for processing during translation. The general translation strategy refers to a balanced optimization translation method for hybrid instruction blocks, which does not favor computation optimization or memory access optimization. Specifically, it includes: using the standard graph coloring algorithm for register allocation during the register allocation stage, using the basic instruction scheduling algorithm to balance the interleaved execution of computation instructions and memory access instructions during the instruction reassembly stage, and using the default thread mapping method during the SIMT mapping stage. For example, for a hybrid instruction block containing 40% computation instructions and 35% memory access instructions, the translation will consider both register utilization and memory access efficiency, and achieve a balance between the two, rather than focusing on loop unrolling as in computationally intensive instruction blocks, or focusing on memory merging optimization as in memory-intensive instruction blocks.

[0051] Then, the priority of each instruction block to be converted is determined. During scheduling and execution, the priority ordering rule is: compute-intensive instruction blocks have the highest priority, mixed instruction blocks have the second highest priority, and memory-intensive instruction blocks have the lowest priority.

[0052] Step 2: Use a lightweight front-end translator to convert the instruction block to be translated into architecture-independent intermediate representation instructions, while preserving the SIMT parallelism and data dependencies of the instruction block to be translated. Then define a translation task, which includes using the target back-end optimizer to sequentially complete instruction selection, register allocation, instruction reassembly, SIMT mapping and instruction encoding of the intermediate representation instructions to obtain the native instructions of the target GPU architecture.

[0053] First, the lightweight front-end translator converts the instruction block to be transformed into architecture-independent intermediate representation instructions, while preserving the parallelism and data dependencies within the instruction block. The transformation steps are as follows:

[0054] Step 2.1: First, perform lexical analysis. Use a lexical analyzer to parse each instruction block to be converted into a structured PTX instruction sequence, and identify the opcode, operands and modifiers of each PTX instruction in the PTX instruction sequence.

[0055] Then, semantic extraction is performed. A semantic analyzer converts each PTX instruction into its corresponding intermediate representation LLVM IR instruction according to predefined mapping rules, resulting in the LLVM IR instruction sequence for each instruction block to be converted. The predefined mapping rules refer to the semantic equivalence conversion rules between PTX instructions and intermediate representation LLVM IR instructions. These rules are predefined and stored in a mapping rule table during the translation framework implementation. The principle of defining the mapping rules is to ensure complete semantic equivalence between the instructions before and after conversion. This ensures that the converted LLVM IR instruction is semantically completely equivalent to the original PTX instruction. For example, the semantics of the PTX multiply-accumulate instruction "mad.f32 %r3, %r1, %r2, %r3" is to multiply %r1 and %r2 and add the original value of %r3, storing the result in %r3. Its corresponding mapping rule is to convert it into the LLVM IR fused multiply-accumulate instruction "%r3 = fma float %r1, %r2, %r3". The mapping rule for the PTX global memory load instruction "ld.global.f32 %r1, [%r2]" is to convert it into the LLVM IR load instruction "%r1 = load float, float* %r2". The mapping rule for the PTX branch instruction "bra target" is to convert it into the LLVM IR unconditional jump instruction "br label %target". These mapping rules cover all instruction types in the PTX instruction set, ensuring that any valid PTX program can be correctly converted into a semantically equivalent LLVM IR representation.

[0056] Step 2.2: Perform SIMT parallelism annotation. Extract the thread block dimension `blockDim` and grid dimension `gridDim` from the kernel startup parameters. Simultaneously, analyze the usage patterns of special registers `%tid`, `%ctaid`, and `%ntid` in the PTX code of the instruction block to be transformed. Append the obtained SIMT parallel structure information as metadata to the generated intermediate representation LLVM IR instructions. The SIMT parallel structure information includes: grid configuration information, i.e., the number of thread blocks in the three dimensions `gridDim.x`, `gridDim.y`, and `gridDim.z`; thread block configuration information, i.e., the number of threads in the three dimensions `blockDim.x`, `blockDim.y`, and `blockDim.z`; thread index calculation pattern, i.e., how special registers are used in the program to calculate the global thread index; shared memory usage information, i.e., the size of shared memory requested by each thread block; and synchronization point information, i.e., the location of synchronization instructions such as `bar.sync` in the program.

[0057] Step 2.3: Perform data dependency analysis, traverse the intermediate representation LLVM IR instruction sequence corresponding to the instruction block to be converted, identify the data dependency relationship between each intermediate representation LLVM IR instruction and other intermediate representation LLVM IR instructions in the sequence, and construct a data flow graph between instructions.

[0058] The data dependencies include RAW dependencies (read-after-write dependencies), WAR dependencies (write-after-read dependencies), WAW dependencies (write-after-write dependencies), and memory dependencies. RAW dependencies (read-after-write dependencies) mean that subsequent instructions need to read the results written by preceding instructions; WAR dependencies (write-after-read dependencies) mean that subsequent instructions will overwrite the data source being read by preceding instructions; WAW dependencies (write-after-write dependencies) mean that multiple instructions write to the same target location sequentially; memory dependencies are the dependencies between video memory accesses identified through pointer alias analysis.

[0059] Then, the translation task is defined, and the target backend optimizer is used to sequentially perform instruction selection, register allocation, instruction reassembly, SIMT mapping, and instruction encoding on the intermediate representation instructions to obtain the native instructions of the target GPU architecture. The specific steps are as follows:

[0060] Step 2.4: Perform instruction selection. According to the instruction set specification of the AMD GCN architecture, map each intermediate LLVMIR instruction to the corresponding AMD GCN architecture instruction to obtain the AMD GCN architecture instruction sequence.

[0061] Step 2.5: Perform register allocation. Since the AMD GCN architecture has strict limitations on the number of vector registers VGPR and scalar registers SGPR, the target back-end optimizer uses a graph coloring algorithm to allocate physical registers for the virtual registers in the intermediate representation of LLVM IR instructions. When the number of physical registers is insufficient, some data is temporarily stored in video memory by inserting register overflow code.

[0062] Step 2.6: Perform instruction reorganization. Based on the scheduling characteristics of Wavefront in the AMD GCN architecture, rearrange the instruction sequence of the AMD GCN architecture. The purpose is to hide memory access latency through instruction-level parallelism, and at the same time, insert s_waitcnt instructions at appropriate positions to manage the synchronization of memory operations.

[0063] Step 2.7: Perform SIMT mapping on the recombined AMD GCN architecture instruction sequence to obtain the target GPU architecture instruction sequence after thread mapping adjustment. The thread organization and parallel execution logic of this target GPU architecture instruction sequence are adapted to the target GPU architecture. In this embodiment, since NVIDIA GPUs use 32 threads as a Warp and AMD GPUs use 64 threads as a Wavefront, the target backend optimizer adjusts the thread index calculation method and branch divergence processing logic of the recombined AMD GCN architecture instruction sequence. The original parallel execution logic based on Warp (32 threads) is adjusted to the execution logic adapted to Wavefront (64 threads). The thread index calculation formula and thread mask operation are adjusted, and NVIDIA's Warp-level primitives (such as __ballot, __shfl, etc.) are converted into AMD's corresponding Wavefront-level operations to obtain the target GPU architecture instruction sequence adapted to the target AMDGCN architecture.

[0064] Step 2.8: Perform instruction encoding. Pack the AMD GCN architecture instruction sequence after SIMT mapping into directly executable machine code according to the binary encoding format of AMD GPU, and obtain the native instruction sequence that can be executed by the target AMD GPU.

[0065] Step 3: Based on the priority order from high to low, allocate AMD GPU thread blocks to the intermediate representation instructions corresponding to the instruction block to be converted according to the preset scheduling strategy. Then, map the batch intermediate representation LLVM IR instruction translation task to the corresponding AMD GPU thread block, use SIMT parallelism for batch processing, translate to obtain the native instructions of the AMD GCN architecture, and store them in the AMD GPU video memory output buffer.

[0066] Specifically, based on the priority determined for each instruction block to be converted in step 1.2, the instruction blocks are sorted in descending order of priority. GPU thread blocks are allocated to the intermediate representation LLVM IR instructions corresponding to each instruction block according to a preset scheduling strategy. Then, SIMT (Single Instruction Multiple Threads) parallelism is used to process the intermediate representation LLVM IR instructions in batches. For example, based on the number of existing GPU thread blocks, every 128 intermediate representation LLVM IR instructions are grouped into a batch, and each batch of intermediate representation LLVM IR instructions is processed in parallel. Each GPU thread block is responsible for translating one or a group of IR instructions. Threads within the same Wavefront / Warp execute the same translation operations (instruction selection, register mapping, encoding) in parallel. Finally, the native instruction sequence executable by the AMD GPU corresponding to each instruction block to be converted is obtained and stored in the GPU's output buffer. This invention uses a parallel translation method to achieve GPU-accelerated translation, reducing the time complexity of traditional CPU serial translation from O(n) to O(n / p), thereby improving translation throughput, where n is the number of intermediate representation LLVM IR instructions and p is the number of parallel GPU thread blocks.

[0067] It should be noted that the default scheduling strategy for allocating GPU thread blocks to the intermediate representation LLVM IR instructions corresponding to the instruction blocks to be translated is as follows: After sorting the instruction blocks to be translated in descending order of priority, when mapping the translation tasks of batch intermediate representation instructions to GPU thread blocks of the AMD GPU, for high-priority computationally intensive instruction blocks, their corresponding intermediate representation instructions are given priority in allocating GPU thread block resources for translation. This ensures that computationally intensive code can be translated and submitted to the AMD GPU for execution as quickly as possible, thereby maximizing the utilization of GPU computing units. When GPU thread block resources are sufficient, mixed instruction blocks and memory-intensive instruction blocks are processed in order of priority. When GPU thread block resources are scarce, low-priority memory-intensive instruction blocks wait for high-priority instruction blocks to be translated before obtaining GPU thread block resources. Through this priority-based default scheduling strategy, computing units are prevented from being idle due to waiting for memory access operations, thus improving the utilization of GPU computing resources during the translation process.

[0068] Furthermore, during the parallel translation process involving multiple GPU thread blocks, each translation stage generates corresponding temporary data and state information, which are referred to as intermediate translation states. The translation stages include five phases: decoding, allocation, scheduling, encoding, and verification. Intermediate translation states include symbol table state, dependency graph state, translation progress state, and intermediate encoding results. Symbol table state refers to the mapping relationship between virtual registers in the original instruction stream and physical registers in the target architecture, as well as the currently allocated and available set of physical registers. Dependency graph state refers to the data dependencies between instructions constructed during the analysis phase of the lightweight front-end translator, represented in the form of an adjacency list, indicating RAW, WAR, WAW, and memory dependencies between instructions. Translation progress state refers to the translation completion flags of each instruction block and the current translation stage. Intermediate encoding results refer to the partially completed native instruction sequence, which may require address backfilling or relocation correction in subsequent stages to correct the symbol table state.

[0069] Since multiple GPU thread blocks may need to access and modify the translation intermediate state simultaneously, this embodiment achieves fast synchronization of the translation intermediate state through shared memory and atomic operations. Specific methods include:

[0070] First, a dedicated area is allocated within the AMD GPU's shared memory to store intermediate translation states, including a symbol table area, a dependency graph area, and a state flag area. Because the access latency of shared memory is lower than that of global memory, state access efficiency can be improved.

[0071] Then, for the first data that only needs to be read, such as the source instruction stream and dependencies, each GPU thread block reads directly from the GPU shared memory in parallel without synchronization. For the second data that needs to be updated by multiple GPU thread blocks, such as register allocation counters and translation completion flags, atomic operations are used to update them to ensure atomicity. For example, the atomicAdd atomic addition operation is used to update the allocated register count, and the atomicCAS atomic comparison and swap operation is used to update the status flags.

[0072] Finally, at the switching points between translation stages, the __syncthreads() barrier instruction is used to synchronize within the thread block, ensuring that all threads in the previous stage have completed their work before entering the next stage, thus avoiding data races and state inconsistencies.

[0073] This embodiment uses shared memory and atomic operations to ensure that all intermediate translation states reside in the GPU's local memory and shared memory. Only after the entire translation process is completed is the final translated binary result's original instruction sequence copied back to CPU memory or executed directly on the GPU. This avoids the latency overhead caused by frequent data transfers between the CPU and GPU via the PCIe bus during the execution of each translation stage, thereby reducing resource consumption.

[0074] Step 4: The AMD GPU executes the translated native instruction sequence that can be executed by the AMD GPU, while collecting performance metrics during the execution of the AMD GPU, and dynamically adjusting the subsequent translation strategy according to the preset feedback mechanism.

[0075] Specifically, the steps for collecting performance metrics during AMD GPU execution and dynamically adjusting subsequent translation strategies based on a preset feedback mechanism include:

[0076] Step 4.1: During the execution of the translated native instruction sequence by the AMD GPU, collect the performance metrics of the AMD GPU through GPU hardware performance counters, including Wavefront pause rate (reflecting the utilization efficiency of computing units), L2 cache miss rate (reflecting memory access locality), branch divergence rate (reflecting SIMT execution efficiency), and register overflow count (reflecting register pressure), etc.

[0077] Step 4.2: Based on the collected performance metrics, analyze the execution bottlenecks and adjust the translation parameters from Step 2 accordingly. These translation parameters include: the scheduling window size during instruction reassembly, the target value for instruction-level parallelism, the interleaving ratio of memory access instructions and computation instructions, etc.; register overflow threshold, overflow code insertion strategy (immediate overflow or delayed overflow), the number of reserved registers, etc.; the aggressiveness of memory merging access, data prefetch distance, and sensitivity to shared memory bank conflict detection, etc.; including branch prediction confidence threshold, predicate execution conversion conditions, branch divergence handling strategy selection, etc.; loop unrolling factor, software pipeline depth, etc. Based on the performance metrics collected in Step 4.1, adjust the translation parameters specifically to improve subsequent translation efficiency. For example, when the wavefront pause rate is too high, increase the instruction scheduling window and the target value for instruction-level parallelism; when the L2 cache miss rate is too high, increase the data prefetch distance and increase the aggressiveness of memory merging access.

[0078] Step 4.3: Finally, for hot code blocks whose performance metrics have not reached the set threshold, return to Step 2 for incremental retranslation, retranslate the code block and replace the original version.

[0079] Step 5: Store the translated native instruction sequence according to the multi-level translation caching strategy, and perform adaptive cache management according to the cache incremental update and eviction strategy.

[0080] Specifically, a multi-level translation caching strategy is employed to store the translated native instruction sequences, avoiding repeated translation of the same instructions and further improving translation efficiency. This multi-level caching strategy includes L1 and L2 caches. The L1 cache is an instruction fragment-level cache, used to store the translation results of frequently occurring basic blocks. It uses the hash value of the instruction block content to be translated as the index key and the corresponding target AMDGCN architecture binary code as the value. The L1 cache capacity is set relatively small (e.g., 64MB) to ensure minimal access latency. The L2 cache is a function-level cache, used to store the translation results of complete kernel functions. It uses the kernel function signature (including the function name and parameter type list) as the index key. The L2 cache capacity is set relatively large (e.g., 512MB) to accommodate more complete function translation results. When the L2 cache is hit, the entire translation process of that kernel function can be completely skipped.

[0081] After capturing the instruction block to be translated in Step 1 (instruction capture phase), the hash value of the instruction block is calculated. Then, the L1 and L2 caches are queried sequentially. If a cache hit occurs in either level, the cached translation result is returned directly, skipping the subsequent Step 2 translation phase, thus significantly reducing response time. If a cache miss occurs, Step 2 is executed normally. After translation, the result is written to the corresponding level of cache: basic block-level translation results are written to the L1 cache, and complete function-level translation results are written to the L2 cache. During the AMD GPU execution of the translated AMD GPU-executable native instruction sequence, the priority of cache entries is updated based on the execution frequency statistics of each code block in the native instruction sequence. Entries with higher execution frequency have higher priority and are less likely to be evicted when cache space is insufficient.

[0082] The incremental cache update strategy is as follows: when the source code of the simulation task program is partially modified, resulting in changes to the content of some instruction blocks to be converted, it is only necessary to re-translate the changed basic blocks and update the corresponding entries in the L1 cache. At the same time, the complete function entries in the L2 cache containing the basic blocks are marked as invalid. The next time they are accessed, a re-translation is triggered. The incremental cache update mechanism avoids the overhead of full re-translation caused by partial modifications.

[0083] The cache eviction strategy is as follows: The LRU (Least Recently Used) algorithm is used as the basic strategy. On this basis, a reinforcement learning model is combined to predict the future reuse probability of each cache entry. The state space of the reinforcement learning model is defined as a tuple of instruction block hash value and historical reuse count. The action space includes two options: retain and evict. The reward function is calculated based on the actual cache hit situation after the decision. The eviction decision is optimized through continuous learning. Entries with high predicted reuse probability are given higher priority in the LRU queue to make them less likely to be evicted. Entries with low predicted reuse probability are evicted first to free up cache space, thereby maximizing the overall hit rate with limited cache capacity.

[0084] The above are merely preferred embodiments of the present invention and are not intended to limit the present invention. Any modifications, equivalent substitutions, and improvements made within the spirit and principles of the present invention should be included within the scope of protection of the present invention.

Claims

1. A method for instruction set conversion in a GPU heterogeneous environment based on simulation software, characterized in that, include: The simulation software uses dynamic instrumentation to intercept the original instruction stream submitted by the simulation task program to the original GPU / CPU for execution, and marks the instruction blocks to be converted. Then, based on the instruction semantic features and hardware resource matching degree, the instruction blocks to be converted are classified and their priorities are determined; A lightweight front-end translator is used to convert the instruction block to be converted into intermediate representation instructions, while preserving the SIMT parallelism and data dependencies of the instruction block to be converted. Then, a translation task is defined, which includes using the target back-end optimizer to sequentially complete instruction selection, register allocation, instruction reassembly, SIMT mapping and instruction encoding of the intermediate representation instructions to obtain the native instructions of the target GPU architecture. According to the priority order from high to low, target GPU thread blocks are allocated to the intermediate representation instructions corresponding to the instruction block to be converted according to the preset scheduling strategy. Then, the translation tasks of batch intermediate representation instructions are mapped to the corresponding target GPU thread blocks, and batch processing is carried out using SIMT parallelism to translate and obtain the native instructions of the target GPU architecture. Among them, methods for using a lightweight front-end translator to convert instruction blocks into intermediate representation instructions while preserving the SIMT parallelism and data dependencies of the instruction blocks to be converted include: Each block of instructions to be converted is parsed into a structured PTX instruction sequence using a lexical analyzer; The semantic analyzer is used to convert each PTX instruction into a corresponding intermediate representation instruction according to a predefined mapping rule, resulting in an intermediate representation instruction sequence for each instruction block to be converted. Analyze and extract the SIMT parallel structure information of the instruction block to be converted, and append the obtained SIMT parallel structure information to the corresponding intermediate representation instruction sequence in the form of metadata; Traverse the intermediate representation instruction sequence corresponding to the instruction block to be converted, identify the data dependency relationship between each intermediate representation instruction and other intermediate representation instructions in the sequence, and construct a data flow graph between instructions.

2. The instruction set conversion method for GPU heterogeneous environments based on simulation software according to claim 1, characterized in that, The method for classifying the instruction blocks to be converted based on instruction semantic features and hardware resource matching degree is as follows: Calculate the ratio of computational instructions to memory access instructions in the instruction block to be converted; If the proportion of computational instructions in the instruction block to be converted exceeds a first preset threshold, the instruction block to be converted will be classified as a computationally intensive instruction block. If the proportion of memory access instructions in the instruction block to be converted exceeds the second preset threshold, the instruction block to be converted will be classified as a memory access intensive instruction block. If the proportion of computational instructions in the instruction block to be converted does not exceed the first preset threshold and the proportion of memory access instructions does not exceed the second preset threshold, then the instruction block to be converted is classified as a mixed instruction block.

3. The instruction set conversion method for GPU heterogeneous environments based on simulation software according to claim 1, characterized in that, The data dependencies include RAW dependencies (read-after-write dependencies), WAR dependencies (read-after-write dependencies), WAW dependencies (write-after-write dependencies), and memory dependencies.

4. The instruction set conversion method for GPU heterogeneous environments based on simulation software according to claim 3, characterized in that, The method for obtaining the native instructions of the target GPU architecture by using the target back-end optimizer to sequentially perform instruction selection, register allocation, instruction reassembly, SIMT mapping, and instruction encoding on the intermediate representation instructions includes: According to the instruction set specification of the target GPU architecture, each intermediate representation instruction is mapped to the corresponding target GPU architecture instruction to obtain the target GPU architecture instruction sequence; A graph coloring algorithm is used to allocate physical registers for the virtual registers in the intermediate representation instructions; Based on the scheduling characteristics of the set of parallel execution threads in the target GPU architecture, the instruction sequence of the target GPU architecture is rearranged, and then SIMT mapping is performed. The SIMT-mapped instruction sequence of the target GPU architecture is packaged into directly executable machine code according to the binary encoding format of the target GPU, thus obtaining the native instruction sequence that can be executed by the target GPU.

5. The instruction set conversion method for GPU heterogeneous environments based on simulation software according to claim 4, characterized in that, When using SIMT parallelism for batch processing, rapid synchronization of intermediate translation states between GPU thread blocks is achieved through shared memory and atomic operations, including the following steps: A dedicated area is allocated in the shared memory of the target GPU to store the intermediate translation state; For the first data that only needs to be read, each GPU thread block reads it in parallel from the GPU shared memory without synchronization between GPU thread blocks; for the second data that needs to be updated by multiple GPU thread blocks, atomic operations are used for updating. At the switching points between translation stages, the intermediate translation state is synchronized between GPU thread blocks using the thread block synchronization barrier instruction.

6. The instruction set conversion method for GPU heterogeneous environments based on simulation software according to claim 5, characterized in that, The intermediate translation state refers to the temporary data and status information generated at each translation stage during the parallel translation process involving multiple GPU thread blocks.

7. The instruction set conversion method for GPU heterogeneous environments based on simulation software according to claim 6, characterized in that, Also includes: The target GPU executes the translated native instructions, collects performance metrics during the execution of the target GPU, and dynamically adjusts the subsequent translation strategy according to a preset feedback mechanism.

8. The instruction set conversion method for GPU heterogeneous environments based on simulation software according to claim 7, characterized in that, Also includes: The translated native instruction sequence is stored according to a multi-level translation caching strategy, and adaptive cache management is performed according to the cache incremental update and eviction strategy.

9. The instruction set conversion method for GPU heterogeneous environments based on simulation software according to claim 8, characterized in that, The multi-level translation caching strategy includes L1 cache and L2 cache; the L1 cache is an instruction fragment level cache, used to store the translation results of frequently occurring basic blocks, with the hash value of the instruction block content to be converted as the index key and the corresponding target GPU architecture binary code as the value; The L2 cache is a function-level cache used to store the translation results of complete kernel functions, with the kernel function signature as the index key.

Images