GPU-oriented compiler constant storage method, apparatus and device, and medium

Abstract

The invention provides a GPU-oriented compiler constant storage method and device, equipment and a medium, and belongs to the field of data processing. According to the method, a control buffer with a dynamic size is divided in a scalar register file and is shared by a whole kernel, a Workload Manager is responsible for space distribution, recovery and initialization, and constants are stored in the control buffer divided by a scalar register during compiling, so that a GPU of an SIMT architecture does not need to have multiple copies in a scalar register file at the same moment, and the number of copies in the scalar register file at the same moment is reduced. And a constant does not need to be obtained from a memory, and the instruction code can accommodate a plurality of 64-bit immediate operands, so that the problem of insufficient resource utilization of a large number of scalar registers in the prior art is solved.

Description

Technical Field

[0001] This invention relates to the field of data processing, and more specifically to a method, apparatus, device, and medium for storing compiler constants for GPUs. Background Technology

[0002] In GPUs with a SIMT (Single Instruction Multiple Thread) architecture, register files are typically implemented using SRAM and are divided into scalar register files and vector register files. The former only supports scalar access and stores uniform data within a single warp, while the latter supports vector access and stores diversity data within a single warp.

[0003] Take double-precision FMA (fused-Multiply-Add) operation as an example.

[0004] smov s[0], 0x3FF80000 smov s[1], 0x00000000 smov s[2], 0x3FE33333 smov s[3], 0x33333333 fma r[0,1], r[2,3], s[2,3], s[0,1] The code snippet above calculates the value of the linear polynomial y = 0.6x + 1.5, a basic numerical operation commonly found in scientific computing programs. The output y is stored in vector registers r0 and r1, and the input x is stored in vector registers r2 and r3. The polynomial coefficients 0.6 (0x3FE3333333333333) and the constant 1.5 (0x3FF8000000000000) are directly encoded in the instruction sequence and loaded into scalar registers s0, s1, s2, and s3 using the smov instruction.

[0005] To improve the utilization of computing units, GPUs typically allow multiple warps to execute simultaneously on a single CU (Compute Unit). For this purpose, a CU needs to be equipped with a large number of scalar registers and must store multiple copies of the same constant at any given time. For example, to support the parallel execution of up to 32 warps, the AMD GCN architecture is equipped with up to 3200 scalar registers per CU, and the same constant may exist in 32 copies simultaneously.

[0006] Furthermore, modern GPU instruction encoding typically cannot accommodate multiple 64-bit immediate values, while FMA instructions often require the use of two constants. There are three common solutions to this problem: first, encode the constants in multiple instructions; second, store the constants in memory; and third, add a constant buffer in the CU at the same level as VL1. The first solution reduces the density of computation instructions, thus reducing the actual peak computing power of the CU; the second solution requires waiting for the VLS unit to asynchronously load the constants into the scalar register, increasing the overall application latency; the third solution requires configuring additional SRAM in the CU and adding complex operand expectation logic, significantly increasing the chip cost. It is clear that each of these solutions has its own problems.

[0007] Therefore, how to improve compilation efficiency and make full use of register resources is an urgent problem to be solved. Summary of the Invention

[0008] The purpose of this invention is to provide a method, apparatus, device, and medium for storing compiler constants for GPUs, so as to at least partially solve the above-mentioned problems.

[0009] According to one aspect of this disclosure, a compiler constant storage method for GPUs is proposed, comprising: Allocate a dynamically sized constant buffer in the scalar register file to store constants during compilation. During assembly, the computed instructions explicitly reference the constant buffer, specifying the number of scalar registers used by a single warp and the capacity of the constant buffer used by the entire kernel. In response to kernel startup, the constant buffer space is allocated and initialized; in response to kernel destruction and reclamation, the constant buffer space is destroyed and reclaimed. In response to warp start, a scalar register is allocated; in response to warp end, the scalar register is reclaimed. The input instruction is decoded, the operand type is identified, and the operand is retrieved according to the operand type.

[0010] In some embodiments, the method further includes that the constant buffer is shared throughout the kernel and that the constant buffer is read-only.

[0011] In some embodiments, the method further includes the Workload Manager being responsible for the allocation, reclamation, and initialization of the constantbuffer space.

[0012] In some embodiments, the method further includes the GPU being a GPU with a single-instruction multithreaded architecture.

[0013] According to another aspect of this disclosure, a compiler constant storage device for GPUs is proposed, comprising: The partitioning module is used to allocate a dynamically sized constant buffer in the scalar register file for storing constants during compilation. Specify the module used during assembly to compute instructions that explicitly reference the constant buffer, specify the number of scalar registers used by a single warp, and specify the capacity of the constant buffer used by the entire kernel. The first response module is used to respond to kernel startup, allocate and initialize constant buffer space, and respond to kernel destruction and reclamation of constant buffer space. The second response module is used to allocate scalar registers in response to warp start and to reclaim scalar registers in response to warp end. The processing module is used to decode the input instructions, identify the type of operands, and retrieve the operands according to their type.

[0014] In some embodiments, the apparatus further includes a constant buffer that is shared throughout the kernel and is read-only.

[0015] In some embodiments, the apparatus further includes a Workload Manager responsible for the allocation, reclamation, and initialization of the constant buffer space.

[0016] In some embodiments, the device further includes a GPU that is a single-instruction multithreaded architecture GPU.

[0017] This application also provides an electronic device, which includes a memory and a processor. The memory stores a computer program, and the processor executes the steps in any of the above embodiments by calling the computer program stored in the memory.

[0018] This application also provides a computer-readable storage medium storing a computer program, characterized in that: the computer program is executed by a processor to perform the steps in any of the above embodiments.

[0019] This invention provides a compiler constant storage method for GPUs. This method is based on allocating a dynamically sized constant buffer in the scalar register file, shared by the entire kernel. The WorkloadManager is responsible for the allocation, reclamation, and initialization of this space. During compilation, constants are stored in the constant buffer allocated from the scalar register. In this way, for GPUs with a SIMT architecture, the same immediate value does not need to exist in multiple copies in the scalar register file at the same time, nor does it need to be retrieved from memory. Instruction encoding can accommodate multiple 64-bit immediate values, solving the problem of insufficient utilization of existing large scalar register resources. Attached Figure Description

[0020] Figure 1 This is a schematic diagram of a compiler constant storage method for GPUs provided in an embodiment of this application.

[0021] Figure 2 This is a schematic diagram of a compiler constant storage device structure for GPUs, provided as an embodiment of this application.

[0022] Figure 3 This is a schematic diagram of the structure of an electronic device provided in an embodiment of this application. Detailed Implementation

[0023] The specific embodiments of the present invention will be further described below with reference to the accompanying drawings. It should be noted that these descriptions are for the purpose of aiding understanding the present invention, but do not constitute a limitation thereof. Furthermore, the technical features involved in the various embodiments of the present invention described below can be combined with each other as long as they do not conflict with each other.

[0024] It should be noted that in the description of this invention, the terms "upper", "lower", "left", "right", "front", "rear", etc., indicate the orientation or positional relationship based on the description of the structure of this invention shown in the accompanying drawings. They are only for the convenience of describing this invention and do not indicate or imply that the device or element referred to must have a specific orientation, or be constructed and operated in a specific orientation. Therefore, they should not be construed as limitations on this invention.

[0025] The terms "first" and "second" in this technical solution are merely designations for corresponding structures that are identical or similar, or that perform similar functions. They do not represent an arrangement of the importance of these structures, nor do they imply any ranking, comparison of size, or other meaning.

[0026] Furthermore, unless otherwise explicitly specified and limited, the terms "installation" and "connection" should be interpreted broadly. For example, a connection can be a fixed connection, a detachable connection, or an integral connection; it can be a mechanical connection or an electrical connection; it can be a direct connection or an indirect connection through an intermediate medium; it can be a connection within two structures. Those skilled in the art can understand the specific meaning of the above terms in this invention by considering the overall concept of the invention and the specific context of the solution.

[0027] This application provides a compiler constant storage method for GPUs. Specifically, please refer to... Figure 1 The basic implementation process of this embodiment is as follows: Step S1: Allocate a dynamically sized constant buffer in the scalar register file to store constants during compilation.

[0028] In a GPU with a SIMT (Single Instruction Multiple Thread) architecture, registers are the fastest component in each SM (Stream Multiprocessor). Register files are typically implemented using SRAM and are divided into scalar register files and vector register files. Scalar register files are used to store scalar data, i.e., single numerical values.

[0029] In GPUs, scalar registers are typically used to store single floating-point numbers, integers, or Boolean values. Scalar register files provide fast access to individual data elements for performing simple arithmetic, logical, or data transfer operations. During compilation, the compiler analyzes the code and determines which scalar registers are needed to store intermediate results, variables, or temporary data. The compiler optimizes register usage to minimize register overflows and memory accesses, thereby improving performance. Vector register files are used to store vector data, which is a collection of data containing multiple values. In GPUs, vector registers are typically used to store collections of multiple floating-point numbers or integers to support vector operations, such as SIMD (Single Instruction Multiple Data) operations. Vector register files can process multiple data elements simultaneously, accelerating mathematical operations and graphics processing. During compilation, the compiler analyzes vector operations in the code and determines which vector registers are needed to store and manipulate data. The compiler optimizes vector register usage to maximize parallelism and throughput, thereby improving GPU performance. Scalar register files only support scalar access and store uniform data within a single warp. Uniform data generally refers to data that will not change due to execution path divergences during parallel processing in SIMD operations. Vector register files support vector access and store diversity data within a single warp. In a SIMD architecture, ideally, all threads execute the same instruction sequence to maximize hardware utilization and performance. However, when divergences exist, the hardware must handle different execution paths of different threads, which typically leads to performance degradation. Divergence data generally refers to the data on the differences in execution paths between parallel threads. In one possible implementation, a dynamically sized constant buffer is allocated within the scalar register file. During compilation or runtime, register space is dynamically allocated to the constant buffer as needed. The size of the constant buffer can be set according to the actual application scenario, and the constant buffer can be shared by the entire kernel.

[0030] Step S2: Calculate the instruction to explicitly reference the constant buffer, specify the number of scalar registers used by a single warp, and specify the capacity of the constant buffer used by the entire kernel.

[0031] In assembly language, designing computation instructions explicitly references a constant buffer, specifying the number of scalar registers used by a single warp. A warp, also known as a thread bundle in a GPU, is a group of consecutive threads, typically consisting of 32 threads (the number can vary depending on the hardware architecture). Threads within a warp execute the same instructions in parallel within the GPU, a model known as SIMD (Single Instruction, Multiple Data). Each SIMD contains tens of thousands of registers, which are allocated to specific threads when the kernel is started. The processor's execution pipeline involves warp switching. The processor's execution pipeline typically consists of five stages: fetch, decode, memory access, execution, and write-back. Memory access here refers to accessing registers. When CUDA is executed, the kernel functions in the host (usually the CPU) are executed on the GPU (device) according to the concept of a thread grid. When the kernel is executed on the GPU, it actually starts many threads. All the threads started by a kernel are called a grid. Threads on the same grid share the same global memory space. In this embodiment, the capacity of the constant buffer is also shared by the entire kernel.

[0032] Step S3: In response to kernel startup, allocate and initialize constant buffer space, and in response to kernel destruction and reclamation of constant buffer space.

[0033] Step S4: In response to warp start, allocate a scalar register; in response to warp end, reclaim the scalar register.

[0034] In one possible implementation, during compilation, the workload manager is responsible for the allocation and reclamation of the entire SRAM; therefore, the constant buffer and the entire scalar register are also managed by the workload manager, including allocation, initialization and destruction / reclamation. When the kernel starts, the constant buffer is allocated and initialized, and reclaimed when destroyed. When the warp starts, the scalar register is allocated and reclaimed when it ends.

[0035] The workload manager is a component or logic responsible for managing and optimizing GPU workloads. Its main task is to receive graphics or computational tasks from applications and break them down into smaller units or work items for parallel execution on the GPU. During compilation, the workload manager can perform the following functions as needed: Task partitioning: Breaking down large graphics rendering or computational tasks into smaller, independently executable tasks. This helps to fully utilize the parallel processing capabilities of the GPU.

[0036] Resource allocation: Allocate appropriate GPU resources, such as compute cores, memory bandwidth, and texture units, to each task based on its requirements. This ensures efficient resource utilization and avoids resource conflicts and waste.

[0037] Optimize scheduling: Based on the dependencies and execution order between work items, optimize their scheduling strategy. This helps reduce latency and improve the overall performance of the GPU.

[0038] Load balancing: If the GPU supports multiple processing units or cores, the workload manager will also be responsible for distributing work items to different cores to achieve load balancing, thereby making full use of all available computing resources.

[0039] Performance Analysis: During the compilation process, the workload manager may also perform performance analysis to identify potential bottlenecks and optimization opportunities. This helps to further improve GPU performance at runtime.

[0040] One possible example of a scalar register file allocation method is shown in Table 1 below.

[0041] Table 1 address use 0x00 warp0 s0 0x04 warp0 s1 0x08 warp1 s0 … … 0x3c warp7 s1 0x40 c0 0x44 c1 Where s0 and s1 represent the scalar register numbers, c0 and c1 represent constant buffers, and warp0 and warp1 are the warp numbers.

[0042] Table 2 shows a possible example of how vector register files can be allocated.

[0043] Table 2

[0044] Where v0 and v1 represent the vector register numbers, warp0 and warp1 are the warp numbers, and thread0 and thread1 are the thread numbers.

[0045] Step S5: Decode the input instruction, identify the type of operand, and retrieve the operand according to the type of operand.

[0046] In this embodiment, the sequencer organizes the input instructions, then decodes them, identifies the operand types, and retrieves the operands based on their types. During compilation, constants are stored in a constant buffer. During GPU compilation, source code (such as code written in high-level shader languages ​​like GLSL or HLSL) is first converted to an intermediate representation (IR), then further optimized and converted into machine code executable by the GPU. In this process, instructions need to be organized into a sequence suitable for GPU execution. The decoder checks the validity of the instructions, ensuring they conform to the GPU's instruction set specification. Next, the GPU identifies the operand types. Operands are the data required to execute instructions; they can be variables, constants, or temporary values ​​stored in GPU memory. The GPU needs to know the operand types to correctly interpret and manipulate this data. For example, if the operand is a floating-point number, the GPU needs to process it according to floating-point arithmetic rules; if the operand is an integer, the arithmetic rules will be different. Finally, the GPU retrieves the operands based on their types. This step involves reading data from the GPU's memory. The GPU will locate the corresponding data in memory based on the address or index of the operand, and load it into the GPU's registers or cache for use when executing instructions later.

[0047] To make the objectives, technical solutions, and advantages of the present invention clearer, the technical solutions of the present invention will be clearly described below in conjunction with a specific embodiment.

[0048] Taking the calculation of the linear function y = 0.6x + 1.5 as an example, 512 threads are enabled on the device side. Assuming each warp contains 64 threads, the Workload Manager will use 8 warps to complete the calculation, as shown in the following code: __global__ void helloFromGPU(double bias, double *out, double *in){ int tid = blockIdx.x * blockDim.x + threadIdx.x; in[tid] = in[tid] * 0.6 + a; } int main() { helloFromGPU<<<1, 512>>>(1.5, out, in); return 0; } After the above code is compiled by the compiler, the device receives the following sequence of assembly instructions: ...... s_ld s[0,1],&bias v_ld v[0,1],&in+tid v_fma v[0,1], v[0,1], c[0,1],s[0,1] v_st &out+tid, v[0,1] ...... .long 858993459 .long 1071854387 Here, v represents the vector register, s represents the scalar register, and c represents the constant buffer. The compiler recognizes that the variable bias is uniform for a single warp, so it loads it into the s register. Similarly, the compiler recognizes that the variable in is diverse for a single warp, so it loads it into the v register. The constant 0.6 is stored in the static constant area of ​​the code.

[0049] When the kernel starts, the Workload Manager allocates a set of independent vector registers for each thread, occupying sizeof(double) × 512 bytes of the vector register file; it allocates a set of independent scalar registers for each warp, occupying sizeof(double) × 8 bytes of the scalar register file; and it allocates a constant buffer for the entire kernel, occupying sizeof(double) × 1 byte of the scalar register file, and sets the initial values ​​according to the static constants in the assembly code. The contents of the constant buffer can only be read by the user and cannot be modified.

[0050] In summary, this invention provides a compiler constant storage method for GPUs. This method is based on allocating a dynamically sized constant buffer in the scalar register file, shared by the entire kernel. The Workload Manager is responsible for the allocation, reclamation, and initialization of this space. During compilation, constants are stored in the constant buffer allocated from the scalar register. In this way, for GPUs with a SIMT architecture, the same immediate value does not need to exist in multiple copies in the scalar register file at the same time, nor does it need to be retrieved from memory. Instruction encoding can accommodate multiple 64-bit immediate values, solving the problem of insufficient utilization of existing large scalar register resources.

[0051] To achieve the above objectives, this embodiment proposes a compiler constant storage device for GPUs. For details, please refer to [link to relevant documentation]. Figure 2 ,include, The partitioning module is used to allocate a dynamically sized constant buffer in the scalar register file for storing constants during compilation. Specify the module used during assembly to compute instructions that explicitly reference the constant buffer, specify the number of scalar registers used by a single warp, and specify the capacity of the constant buffer used by the entire kernel. The first response module is used to respond to kernel startup, allocate and initialize constant buffer space, and respond to kernel destruction and reclamation of constant buffer space. The second response module is used to allocate scalar registers in response to warp start and to reclaim scalar registers in response to warp end. The processing module is used to decode the input instructions, identify the type of operands, and retrieve the operands according to their type.

[0052] In some embodiments, the apparatus further includes a constant buffer that is shared throughout the kernel and is read-only.

[0053] In some embodiments, the apparatus further includes a Workload Manager responsible for the allocation, reclamation, and initialization of the constant buffer space.

[0054] In some embodiments, the device further includes a GPU that is a single-instruction multithreaded architecture GPU.

[0055] Accordingly, embodiments of this application also provide an electronic device, which can be a terminal or a server. For example... Figure 3 As shown, Figure 3 This is a schematic diagram of the structure of an electronic device provided in an embodiment of this application.

[0056] The electronic device 300 includes a processor 301 with one or more processing cores, a memory 302 with one or more computer-readable storage media, and a computer program stored in the memory 302 and executable on the processor. The processor 301 and the memory 302 are electrically connected. Those skilled in the art will understand that the electronic device structure shown in the figures does not constitute a limitation on the electronic device, and may include more or fewer components than shown, or combine certain components, or have different component arrangements.

[0057] The processor 301 is the control center of the electronic device 300. It connects various parts of the electronic device 300 through various interfaces and lines. By running or loading software programs (computer programs) and / or units stored in the memory 302, and calling data stored in the memory 302, it performs various functions of the electronic device 300 and processes data, thereby monitoring the electronic device 300 as a whole.

[0058] In this embodiment, the processor 301 in the electronic device 300 loads the instructions corresponding to the processes of one or more applications into the memory 302 according to the following steps, and the processor 301 runs the applications stored in the memory 302 to realize various functions: Allocate a dynamically sized constant buffer in the scalar register file to store constants during compilation. During assembly, the computed instructions explicitly reference the constant buffer, specifying the number of scalar registers used by a single warp and the capacity of the constant buffer used by the entire kernel. In response to kernel startup, the constant buffer space is allocated and initialized; in response to kernel destruction and reclamation, the constant buffer space is destroyed and reclaimed. In response to warp start, a scalar register is allocated; in response to warp end, the scalar register is reclaimed. The input instruction is decoded, the operand type is identified, and the operand is retrieved according to the operand type.

[0059] For details on the specific implementation of each of the above operations, please refer to the aforementioned embodiments, which will not be repeated here.

[0060] Optional, such as Figure 3 As shown, the electronic device 300 further includes: a constant storage module 303, a communication module 304, an input unit 305, and a power supply 306. The processor 301 is electrically connected to the interrupt reporting module 303, the communication module 304, the input unit 305, and the power supply 306. Those skilled in the art will understand that... Figure 3 The electronic device structure shown does not constitute a limitation on the electronic device and may include more or fewer components than shown, or combine certain components, or have different component arrangements.

[0061] The constant storage module 303 can be used to implement compiler constant storage for GPUs.

[0062] The communication module 304 can be used to communicate with other devices.

[0063] The input unit 305 can be used to receive input numbers, characters, or user characteristic information (such as fingerprints, iris, facial information, etc.), and to generate keyboard, mouse, joystick, optical, or trackball signal inputs related to user settings and function control.

[0064] Power supply 306 is used to supply power to various components of electronic device 300. Optionally, power supply 306 can be logically connected to processor 301 through a power management system, thereby enabling functions such as charging, discharging, and power consumption management through the power management system. Power supply 306 may also include one or more DC or AC power supplies, recharging systems, power fault detection circuits, power converters or inverters, power status indicators, and other arbitrary components.

[0065] In the above embodiments, the descriptions of each embodiment have different focuses. For parts not described in detail in a certain embodiment, please refer to the relevant descriptions in other embodiments.

[0066] Those skilled in the art will understand that all or part of the steps in the various methods of the above embodiments can be performed by instructions, or by instructions controlling related hardware. These instructions can be stored in a computer-readable storage medium and loaded and executed by a processor.

[0067] Therefore, embodiments of this application provide a computer-readable storage medium storing multiple computer programs that can be loaded by a processor to execute the steps of a GPU-oriented compiler constant storage method provided in this application. For example, the computer program can execute the following steps: Allocate a dynamically sized constant buffer in the scalar register file to store constants during compilation. During assembly, the computed instructions explicitly reference the constant buffer, specifying the number of scalar registers used by a single warp and the capacity of the constant buffer used by the entire kernel. In response to kernel startup, the constant buffer space is allocated and initialized; in response to kernel destruction and reclamation, the constant buffer space is destroyed and reclaimed. In response to warp start, a scalar register is allocated; in response to warp end, the scalar register is reclaimed. The input instruction is decoded, the operand type is identified, and the operand is retrieved according to the operand type.

[0068] For details on the implementation of each of the above operations, please refer to the previous examples, which will not be repeated here.

[0069] The computer-readable storage medium may include: read-only memory (ROM), random access memory (RAM), disk or optical disk, etc.

[0070] Since the computer program stored in the storage medium can execute the steps in any of the GPU-oriented compiler constant storage methods provided in the embodiments of this application, the beneficial effects that any of the GPU-oriented compiler constant storage methods provided in the embodiments of this application can achieve can be realized. For details, please refer to the previous embodiments, which will not be repeated here.

[0071] This application is described with reference to flowchart illustrations and / or block diagrams of methods, apparatus (systems), and computer program products according to embodiments of this application. It will be understood that each block of the flowchart illustrations and / or block diagrams, and combinations of blocks in the flowchart illustrations and / or block diagrams, can be implemented by computer program instructions. These computer program instructions can be provided to a processor of a general-purpose computer, special-purpose computer, embedded processor, or other programmable data processing apparatus to produce a machine, such that the instructions, which execute via the processor of the computer or other programmable data processing apparatus, generate instructions for implementing the flowchart... Figure 1 One or more processes and / or boxes Figure 1 A device that provides the functions specified in one or more boxes.

[0072] These computer program instructions may also be stored in a computer-readable storage medium that can direct a computer or other programmable data processing device to function in a particular manner, such that the instructions stored in the computer-readable storage medium produce an article of manufacture including instruction means, which are implemented in a process Figure 1 One or more processes and / or boxes Figure 1 The function specified in one or more boxes.

[0073] These computer program instructions may also be loaded onto a computer or other programmable data processing equipment to cause a series of operational steps to be performed on the computer or other programmable equipment to produce a computer-implemented process, thereby providing instructions that execute on the computer or other programmable equipment for implementing the process. Figure 1 One or more processes and / or boxes Figure 1 The steps of the function specified in one or more boxes.

[0074] The embodiments of the present invention have been described in detail above with reference to the accompanying drawings, but the present invention is not limited to the described embodiments. For those skilled in the art, various changes, modifications, substitutions, and variations can be made to these embodiments without departing from the principles and spirit of the present invention, and these variations still fall within the protection scope of the present invention.

Claims

1. A compiler constant storage method for GPUs, characterized in that, The method includes: Allocate a dynamically sized constant buffer in the scalar register file to store constants during compilation. During assembly, the computed instructions explicitly reference the constant buffer, specifying the number of scalar registers used by a single warp and the capacity of the constant buffer used by the entire kernel. In response to kernel startup, the constant buffer space is allocated and initialized; in response to kernel destruction and reclamation, the constant buffer space is destroyed and reclaimed. In response to warp start, a scalar register is allocated; in response to warp end, the scalar register is reclaimed. The input instruction is decoded, the operand type is identified, and the operand is retrieved according to the operand type.

2. The method according to claim 1, characterized in that: The constant buffer is shared by the entire kernel and is read-only.

3. The method according to claim 1, characterized in that: The Workload Manager is responsible for the allocation, reclamation, and initialization of the constant buffer space.

4. The method according to any one of claims 1-3, characterized in that: The GPU in question is a single-instruction, multi-threaded GPU.

5. A compiler constant storage device for GPUs, characterized in that, The device includes: The partitioning module is used to allocate a dynamically sized constant buffer in the scalar register file for storing constants during compilation. Specify the module used during assembly to compute instructions that explicitly reference the constant buffer, specify the number of scalar registers used by a single warp, and specify the capacity of the constant buffer used by the entire kernel. The first response module is used to respond to kernel startup, allocate and initialize constant buffer space, and respond to kernel destruction and reclamation of constant buffer space. The second response module is used to allocate scalar registers in response to warp start and to reclaim scalar registers in response to warp end. The processing module is used to decode the input instructions, identify the type of operands, and retrieve the operands according to their type.

6. The apparatus according to claim 5, characterized in that, The device further includes: The constant buffer is shared by the entire kernel and is read-only.

7. The apparatus according to claim 6, characterized in that, The device further includes: The Workload Manager is responsible for the allocation, reclamation, and initialization of the constant buffer space.

8. The apparatus according to any one of claims 5-7, characterized in that: The GPU in question is a single-instruction, multi-threaded GPU.

9. An electronic device, characterized in that: The method includes a memory storing executable program code and a processor coupled to the memory; wherein the processor invokes the executable program code stored in the memory to perform the method as described in any one of claims 1-4.

10. A computer-readable storage medium storing a computer program, characterized in that: The computer program is executed by the processor to perform the method as described in any one of claims 1-4.

Images