Scalar Access System and Method Based on Vector Memory

Abstract

This invention discloses a scalar access system and method based on vector memory. The system includes a data cache module, a wavefront instruction slot module, a scalar scheduler, and a vector storage module. The data cache module queries the tag block of a received scalar command to see if a tag block with the same target access address as the scalar command exists. If not, it determines the target tag block based on the metadata in the tag block and generates corresponding request information based on the operation type of the scalar command. The scalar scheduler queries the operation type in the request information and accesses its write request buffer or vector storage module based on the operation type to obtain return information. The data cache module performs a storage access operation on the target tag block corresponding to the return information to obtain a first operation result. The first operation result is then transmitted to the wavefront instruction slot module. This invention achieves efficient storage access scheduling and instruction slot release management, improving the overall performance and system throughput of scalar command execution.

Description

Technical Field

[0001] This invention relates to the field of parallel processing technology, and in particular to a scalar access system and method based on vector memory. Background Technology

[0002] In existing parallel processors, computational tasks are typically executed in waves. Each wave consists of multiple threads, and in graphics processing units (GPUs) or general-purpose parallel architectures, multiple waves can execute in parallel. However, the execution order between different waves is often unpredictable, which differs significantly from the strict instruction-order serial execution mode of a central processing unit (CPU). When there are dependencies between multiple waves, synchronization barriers need to be inserted into the computational kernel to specify the execution order, ensuring that subsequent waves can continue running after the first wave has completed. The number, functional characteristics, and access latency of synchronization barrier resources directly affect the execution efficiency of the computational kernel, and thus the overall computational performance of the parallel processor.

[0003] Vector memory (VM) is a key data support unit in parallel processors. It is usually configured with multiple independent dual-port memory units, which have a large storage capacity and can respond to multiple access requests at the same time. This enables efficient vector read and write operations of multiple Single Instruction Multiple Data (SIMD) modules in the Compute Unit, thus supporting high-concurrency computing tasks.

[0004] In existing technologies, synchronization barrier operations are mainly implemented in two ways: First, through special synchronization instructions supported by hardware, such as the barrier instruction used in Compute Unified Device Architecture (CUDA). This method has the advantages of fast response speed and ease of use, but is limited by hardware resource constraints, resulting in a limited number of synchronization barriers that can be supported. Second, through performing atomic operations on specific memory addresses, such as reserving specific addresses in memory space and implementing synchronization control through atomic addition and subtraction. Although this method can support more synchronization resources, the target memory area is often far from the computing unit, resulting in relatively high access latency and thus lower overall synchronization efficiency than the special instruction method. It is usually used as a supplement to the first method.

[0005] Currently, in order to compensate for the shortcomings of the above two synchronization methods in terms of resource quantity or access efficiency, those skilled in the art have improved the flexibility and execution efficiency of synchronization operations by directly configuring synchronization resources in the vector storage module and introducing mechanisms such as multi-thread barrier instruction (mbarrier).

[0006] However, since the vector storage module is essentially a high-speed channel providing vector data services to the processing element (PE), its architecture is primarily designed for high-concurrency vector parallel read / write optimization and typically lacks the ability to directly handle scalar access requests or atomic operations. When the control queue (CQ) needs to access barrier resources in the vector storage module, it often relies on the processing element to complete this process, indirectly calling computational resources such as the Vector General Purpose Register (VGPR) and the Arithmetic Logic Unit (ALU), resulting in additional consumption of computational resources and reduced overall utilization. More importantly, barrier operations are usually single-threaded control logic, which contradicts the internal architecture of the vector storage module, which is optimized for multi-threaded concurrent services. This easily leads to a waste of bandwidth resources in the vector storage module, ultimately affecting the service efficiency of the vector storage module and the overall execution performance of the parallel processor.

[0007] This section is intended to provide background or context for the embodiments of the invention set forth in the claims. The description herein is not an admission that it is prior art simply because it is included in this section. Summary of the Invention

[0008] This invention provides a scalar access system based on vector memory, which optimizes data access paths according to different operation types during scalar command execution, achieving efficient memory access scheduling and instruction slot release management, and improving the overall performance and system throughput of scalar command execution.

[0009] The scalar access system based on vector memory includes: a data cache module, a wavefront instruction slot module, a scalar scheduler, and a vector storage module;

[0010] The data caching module is used to query whether there is a tag block in the tag block of the received scalar command that has the same target access address as the scalar command; if not, it determines the target tag block based on the metadata in the tag block, and generates corresponding request information based on the operation type of the scalar command.

[0011] The scalar scheduler includes a write request buffer. The scalar scheduler is used to query the operation type in the request information, access the write request buffer or the vector storage module based on the operation type, obtain return information, and send the return information to the data cache module through the first data transmission channel.

[0012] The data caching module is used to perform a storage access operation on the target tag block corresponding to the returned information to obtain a first operation result; and transmits the first operation result to the wavefront instruction slot module through the second data transmission channel.

[0013] In some embodiments, the return information includes first return information and second return information; if the operation type is a read operation, the scalar scheduler determines whether the target access address in the request information exists in its write request buffer; if the target access address exists, a data forwarding operation is performed to obtain the first return information; if the target access address does not exist, the request information is stored in the write request buffer; the request information is sent to the vector storage module based on the storage access cycle permission to obtain the second return information.

[0014] In some embodiments, the returned information includes the second returned information. If the operation type is a write operation, the scalar scheduler stores the request information in the write request buffer; and sends the request information to the vector storage module based on the storage access cycle permission to obtain the second returned information.

[0015] In some embodiments, the vector storage module further includes multiple storage arrays; when the scalar scheduler has storage access cycle permissions, the scalar scheduler is also used to query the write request buffer according to the target access address of the request information, determine multiple request information that do not have storage array access conflicts, and send them in parallel to the multiple storage arrays of the vector storage module to obtain the second return information.

[0016] In some embodiments, the system further includes: a vector scheduler; when the scalar scheduler does not have storage access cycle permissions, the scalar scheduler is also used to monitor the occupancy status of each storage array in the vector storage module; if a storage array in an idle state is detected, and the operation type in the request information is the same as the operation type of the storage access operation of the vector scheduler, the request information is inserted into the storage access cycle and sent in parallel to the corresponding storage array to obtain the second return information.

[0017] In some embodiments, the data caching module includes:

[0018] The tag block determination unit is used to determine the target tag block based on the metadata in the tag block according to a first preset rule;

[0019] The command processing unit is used to perform metadata update operations on the target tag block according to the operation type in the scalar command and a second preset rule. When the operation type is a read operation or an atomic operation, the unit generates corresponding request information and sends it to the scalar scheduler.

[0020] In some embodiments, the metadata includes: a valid flag bit; the tag block determination unit is specifically used to query the valid flag bit of the tag block, and if there is a tag block in an idle state, select the tag block with the longest idle time as the target tag block; if there is no tag block in an idle state, select the target tag block using the least recently used algorithm, and generate the request information based on the metadata in the target tag block, and send it to the scalar scheduler through the third data transmission channel.

[0021] In some embodiments, the operation type is a write operation; the command processing unit is specifically used to write the metadata of the scalar command into the target tag block.

[0022] In some embodiments, the operation type is a read operation; the command processing unit is specifically used to set the in-flight flag bit in the target tag block to 1, write the meta-information in the scalar command into the target tag block, generate the request information, and send the request information to the scalar scheduler through the third data transmission channel.

[0023] In some embodiments, the operation type is an atomic operation; the command processing unit is specifically used to set the in-flight flag bit in the target tag block to 1, write the operand in the scalar command into the data field of the target tag block, generate the corresponding request information, and send the request information to the scalar scheduler through the third data transmission channel.

[0024] In some embodiments, the wavefront instruction slot module includes multiple wavefront units, each wavefront unit including multiple instruction slots; the wavefront instruction slot module is used to receive the scalar command, query the occupancy status of all instruction slots in the wavefront unit corresponding to the wavefront identifier based on the wavefront identifier in the scalar command; if there is an idle instruction slot, the scalar command is stored in that instruction slot.

[0025] In some embodiments, if a tag block exists that is the same as the target access address of the scalar command, the data caching module performs a storage access operation on the tag block to obtain a second operation result; wherein, the storage access operation includes: a read operation, a write operation, and an atomic operation.

[0026] In some embodiments, the wavefront command slot module is further configured to determine the corresponding command slot based on the wavefront identifier and slot identifier in the first operation result and the second operation result, and release the command slot according to a third preset rule.

[0027] This invention also provides a scalar access method based on vector memory, applied to the aforementioned scalar access system based on vector memory. This method optimizes data access paths according to different operation types during scalar command execution, achieving efficient memory access scheduling and instruction slot release management, thereby improving the overall performance of scalar command execution and system throughput.

[0028] This vector memory-based scalar access method includes:

[0029] The data caching module queries the received scalar command to see if there is a tag block in its tag block that has the same target access address as the scalar command; if not, it determines the target tag block based on the metadata in the tag block and generates corresponding request information based on the operation type of the scalar command.

[0030] The scalar scheduler includes a write request buffer. The scalar scheduler queries the operation type in the request information, accesses the write request buffer or the vector storage module based on the operation type, obtains return information, and sends the return information to the data cache module through the first data transmission channel.

[0031] The data caching module performs a storage access operation on the target tag block corresponding to the returned information to obtain a first operation result; the first operation result is then transmitted to the wavefront instruction slot module through the second data transmission channel.

[0032] This invention also provides a computer device, including a memory, a processor, and a computer program stored in the memory and executable on the processor. When the processor executes the computer program, it implements the above-described scalar access method based on vector memory.

[0033] This invention also provides a computer-readable storage medium storing a computer program that, when executed by a processor, implements the above-described scalar access method based on vector memory.

[0034] This invention also provides a computer program product, which includes a computer program that, when executed by a processor, implements the above-described scalar access method based on vector memory.

[0035] The scalar access system and method based on vector memory provided in this invention achieves efficient access and data return of scalar commands in the event of a cache miss by introducing a collaborative processing mechanism of a data caching module, a scalar scheduler, and a vector storage module into the scalar access system. The data caching module quickly locates the target tag block based on the tag block's metadata and generates accurate request information based on the operation type, reducing invalid accesses. The scalar scheduler flexibly selects to access the write request buffer or directly access the vector storage module according to the operation type, and can return results immediately when data is available, shortening access latency. Furthermore, the separate design of the first and second data transmission channels makes the data return and result write-back paths independent, reducing channel contention and improving instruction processing parallelism and system throughput. The above scheme improves scalar access efficiency and storage resource utilization while ensuring data consistency. Attached Figure Description

[0036] To more clearly illustrate the technical solutions in the embodiments of the present invention or the prior art, the drawings used in the description of the embodiments or the prior art will be briefly introduced below. Obviously, the drawings described below are only some embodiments of the present invention. For those skilled in the art, other drawings can be obtained based on these drawings without creative effort. In the drawings:

[0037] Figure 1 This is a schematic diagram of a scalar access system based on a vector memory according to an embodiment of the present invention;

[0038] Figure 2 This is a schematic diagram of a scalar access system based on a vector memory in another embodiment of the present invention;

[0039] Figure 3 This is a schematic diagram of the physical structure of an electronic device provided in an embodiment of the present invention. Detailed Implementation

[0040] To make the objectives, technical solutions, and advantages of the embodiments of the present invention clearer, the embodiments of the present invention will be further described in detail below with reference to the accompanying drawings. Here, the illustrative embodiments of the present invention and their descriptions are used to explain the present invention, but are not intended to limit the present invention.

[0041] To effectively support high-speed access and data interaction of scalar commands within a vector memory architecture, and to address issues such as lack of flexible scheduling, low data hit rate, high memory access latency, and imperfect forwarding mechanisms in existing technologies, this application provides a scalar access system based on vector memory. This system uses a data cache module to query and determine the hit rate of the target access address corresponding to the scalar command. If a hit occurs, a target tag block is selected based on metadata, and request information is generated. The scalar scheduler selects to access the write request buffer or the vector memory (VM) module based on the operation type and obtains the corresponding return information. This information is then returned to the data cache module via a first data transmission channel to complete the storage access of the target tag block, forming the operation result. Finally, the result is transmitted to the wavefront instruction slot module via a second data transmission channel. This achieves efficient scalar request hit rate and data forwarding and parallel access, improving the interaction efficiency between scalar commands and vector memory units and reducing memory access latency.

[0042] In one embodiment of this application, a scalar access system based on a vector memory is provided.

[0043] like Figure 1 As shown, the scalar access system includes: a data cache module, a wave instruction slot module (wave_inst_slot), a scalar scheduler (s_sched), and a vector memory bank module (VM).

[0044] The data caching module is used to query the tag of a received scalar command (scmd) to see if there is a tag with the same target access address as the scalar command. If not, the target tag is determined based on the metadata in the tag, and the corresponding request information is generated based on the operation type of the scalar command.

[0045] The scalar scheduler includes a write request buffer. The scalar scheduler queries the operation type in the request information, accesses the write request buffer or vector storage module based on the operation type, and obtains the return information. This return information is then sent to the data cache module via the first data transmission channel (d channel).

[0046] The data caching module performs storage access operations on the target tag block corresponding to the returned information to obtain the first operation result. The first operation result is then transmitted to the wavefront command slot module through the second data transmission channel (e-channel).

[0047] According to the above embodiments, by introducing a collaborative processing mechanism of a data caching module, a scalar scheduler, and a vector storage module into the scalar access system, efficient access and data return of scalar commands in the event of a cache miss are achieved. The data caching module quickly locates the target tag block based on the tag block's metadata and generates accurate request information based on the operation type, reducing invalid accesses. The scalar scheduler flexibly selects to access the write request buffer or directly access the vector storage module according to the operation type, enabling immediate return of results when data is available, shortening access latency. Furthermore, the separate design of the first and second data transmission channels makes the paths for data return and result write-back independent, reducing channel contention and improving instruction processing parallelism and system throughput. The above scheme improves scalar access efficiency and storage resource utilization while ensuring data consistency.

[0048] In embodiments of the present invention, such as Figure 1 As shown, the data caching module contains 32 tag blocks, each storing metadata to describe its status information. This metadata includes fields such as: valid flag (Tag_valid), data sign (Data_sign), data field (Data), data type (Data_type), operation type (Opcode), in-flight flag (Inflight), LRU ord (Lru_ord), wavefront identifier (Wave_id), slot identifier (Slot_id), and target access address (Addr).

[0049] The detailed description of the metadata is shown in Table 1 below:

[0050]

[0051] Table 1

[0052] In embodiments of the present invention, such as Figure 1As shown, when a scalar command sent by the Control Queue (CQ) arrives at the data caching module, the data caching module searches its internal 32 tag blocks for a tag block whose target access address matches the access address of the scalar command. The metadata in the scalar command includes: operation type (Opcode, e.g., atomic_add), target access address, operand, wavefront identifier, and slot identifier. If no tag block exists with the same target access address as the scalar command, the target tag block is determined based on the metadata in the tag block, and corresponding request information is generated based on the operation type of the scalar command. This request information is then sent to the scalar scheduler via channel a. This request information includes read requests and write requests, with corresponding read and write operations, respectively.

[0053] If the operation type (Opcode) of the scalar command is a write operation, the scalar command can directly write the write data to the target tag block in the data cache module after selecting the target tag block, without accessing the vector storage module and receiving return information.

[0054] like Figure 1 As shown, the scalar scheduler includes a write request buffer (wr_req) and a return data array. The write request buffer stores request information from the data cache module sequentially for forwarding operations when necessary. The return data array stores data returned by the vector storage module or data obtained directly through the forwarding mechanism, and returns this data sequentially to the data cache module via the d-channel.

[0055] For scalar commands with a read operation type, after the scalar scheduler completes data reading from the vector storage module or write request buffer, it transmits the read data and its related metadata as return information to the return data array via channel C. The return data array stores this return information in the data cache module. The data cache module determines the corresponding target tag block based on the tag block identifier (tag_id) carried in the return information and updates the metadata stored in the target tag block according to the return information. After the data cache module completes the update, it constructs scalar completion information as the first operation result of the read operation and returns it to the wavefront command slot module via channel E to notify the wavefront command slot module that the scalar command execution is complete.

[0056] For scalar commands with atomic operation types, after the scalar scheduler completes reading the target data from the vector storage module or write request buffer, it transmits the read data and its related metadata as second return information to the return data array via the C channel. The return data array stores this return information in the data cache module. The data cache module determines the corresponding target tag block based on the tag block identifier carried in the return information and performs the corresponding atomic operation (such as atomic addition or atomic subtraction) based on the return information. After the atomic operation is completed, scalar completion information is generated and sent to the wavefront command slot module via the E channel. The scalar completion information is constructed as the first operation result of the atomic operation and returned to the wavefront command slot module via the E channel to notify the wavefront command slot module that the scalar command has been executed. The first operation result includes information such as the wavefront identifier, slot identifier, and tag block identifier.

[0057] In some embodiments, if a tag block exists that has the same target access address as the scalar command, the data caching module performs a storage access operation on the tag block to obtain a second operation result. The storage access operation includes: a read operation, a write operation, and an atomic operation.

[0058] In this embodiment of the invention, if a tag block exists in the data caching module that matches the target access address of a scalar command, the scalar command access is determined to be hit. The data caching module can directly perform the corresponding storage access operation on the hit tag block based on the operation type in the scalar command. The storage access operation includes read operations, write operations, and atomic operations (including addition and subtraction operations).

[0059] After completing the above storage access operation, a scalar completion message is constructed as the second operation result of the read operation and returned to the wavefront command slot module through the e channel to notify the wavefront command slot module that the scalar command has been executed.

[0060] The second operation result includes information such as wavefront identifier, slot identifier, and tag block identifier.

[0061] According to the above embodiments, when the data cache module is hit, the technical solution can directly complete the storage access operation of the target tag block and return the second result containing the necessary information to the wavefront instruction slot module, thereby effectively reducing memory access latency and improving the overall efficiency of instruction scheduling and execution.

[0062] In some embodiments, such as Figure 1 As shown, the wavefront instruction slot module includes multiple wavefront units (e.g., w0, w1, ..., w39), and each wavefront unit includes multiple instruction slots (e.g., ...). Figure 1(As shown by the gray square). The wavefront command slot module is used to receive scalar commands and query the occupancy status of all command slots in the wavefront unit corresponding to the wavefront identifier based on the wavefront identifier in the scalar command. If there is an idle command slot, the scalar command is stored in that command slot.

[0063] In embodiments of the present invention, such as Figure 1 As shown, the wavefront instruction slot module manages the execution order of multiple scalar instructions within each wavefront unit (Wave). This wavefront instruction slot module includes multiple wavefront units (e.g., w0, w1, ..., w39), each with a set of instruction slots below it.

[0064] For example, in this embodiment, each wavefront unit includes 16 instruction slots (e.g., slot0, slot1, ..., slot16), where w0, w1, ..., w39 are wavefront identifiers, and slot0, slot1, ..., slot16 are slot identifiers.

[0065] When a scalar command is input to the wavefront instruction slot module via the data transmission channel, the wavefront instruction slot module determines the target storage location based on the wavefront identifier in the scalar command. Specifically, it first locates the corresponding wavefront cell row area based on the wavefront identifier, and then traverses the occupancy flag bits corresponding to the slot identifiers of all instruction slots in that row area.

[0066] If the occupancy flag corresponding to the slot identifier is 0, it indicates that the instruction slot is in an idle state. In this case, the scalar command is written to the instruction slot, and the metadata field of the scalar command is stored. The metadata field includes: wavefront identifier, slot identifier, opcode, target access address, and data type. The occupancy flag can be encoded in binary or multi-state encoding; this invention does not limit this.

[0067] If the occupancy flag corresponding to the slot identifier is 1, it means that the instruction slot is occupied. Then continue to traverse and search for the next instruction slot until an empty instruction slot is found and the scalar command metadata field is written.

[0068] According to the above embodiments, by setting multiple instruction slots for each wavefront unit in the wavefront instruction slot module, and combining the wavefront identifier and slot identifier occupancy flag management mechanism, precise positioning and orderly storage of scalar commands are achieved, thereby ensuring the sequential consistency and rational resource allocation of multiple scalar instructions within the same wavefront during execution. By dynamically allocating idle instruction slots through traversing the occupancy flags, command waiting or command loss problems caused by slot conflicts can be effectively avoided, improving the real-time performance of command reception and command scheduling.

[0069] In some embodiments, such as Figure 1 As shown, the data caching module includes:

[0070] The tag block determination unit is used to determine the target tag block based on the metadata in the tag block and according to a first preset rule.

[0071] The command processing unit is used to perform metadata update operations on the target tag block according to the operation type in the scalar command and a second preset rule. When the operation type is a read operation or an atomic operation, it generates corresponding request information and sends it to the scalar scheduler. The tag block determination unit and the command processing unit are virtual logical units.

[0072] In this embodiment of the invention, if no tag block matching the target access address of a scalar command exists in the data caching module, the scalar command access is determined to be a miss. The data caching module first determines the target tag block based on the metadata in the tag block according to a first preset rule; then, based on the operation type in the scalar command, it performs a metadata update operation on the target tag block according to a second preset rule. When the operation type is a read operation or an atomic operation, corresponding request information is generated and sent to the scalar scheduler to trigger subsequent memory access processing.

[0073] According to the above embodiments, the data caching module can quickly and accurately locate target tag blocks based on the tag block's metadata, reducing latency caused by invalid searches. The data caching module flexibly adjusts its processing strategy according to different operation types. For write operations, it can directly update metadata internally; for read or atomic operations, it generates request information and sends it to the scalar scheduler, thereby achieving ordered scheduling of memory access requests. The data caching module effectively improves the efficiency of tag block lookup and command processing, reduces storage access waiting time, and by generating request information and sending it to the scalar scheduler, reduces the impact on the backend storage system, ensuring high-concurrency processing capabilities for scalar access and the overall stability of the system.

[0074] In some embodiments, the metadata includes a valid flag. The tag block determination unit is specifically configured to: query the valid flag of a tag block; if an idle tag block exists, select the tag block with the longest idle time as the target tag block; if no idle tag block exists, select the target tag block using the least recently used algorithm, generate request information based on the metadata in the target tag block, and send it to the scalar scheduler through the third data transmission channel (channel a).

[0075] In this embodiment of the invention, if no tag block matching the target access address of the scalar command exists in the data cache module, the valid flag bits of all tag blocks in the data cache module are first traversed. When a tag block is detected to have a valid flag bit of 0 (Tag_valid=0), it indicates that the tag block is in an idle state. If only one idle tag block exists, it is directly used as the target tag block. If multiple idle tag blocks exist, their LRU flags are compared, and the tag block with the largest flag value is selected as the target tag block according to the Least Recently Used (LRU) strategy. After selecting the target tag block, the metadata of the scalar command (including wavefront identifier, slot identifier, operation type, target access address, and data type, etc.) is written into the target tag block, and the corresponding request information is generated according to the operation type of the scalar command and sent to the scalar scheduler through channel a.

[0076] If no free tag block exists in the data cache module (Tag_valid=1), the tag block with the largest value is selected from all tag blocks according to the LRU policy for replacement. The original data in this tag block is packaged as write data, and a write request is generated as the request information. The request information and the write data are sent to the scalar scheduler through channel a.

[0077] After completing the above allocation operation, the data caching module updates the LRU position of the tag block according to the LRU policy to maintain the timeliness of the replacement policy.

[0078] According to the above embodiments, by prioritizing the use of free tag blocks in the event of a cache miss, introducing an LRU-based selection strategy when multiple free tag blocks exist, and replacing tags according to the LRU strategy when no free tag blocks exist, efficient allocation and utilization of data cache resources are achieved. This mechanism reduces cache conflicts while ensuring a high hit rate, preventing the accidental deletion of high-frequency data due to random replacement, thereby reducing memory access latency. Furthermore, by generating and sending request information to the scalar scheduler simultaneously with writing the target tag block, parallel processing of instruction scheduling and cache updates can be achieved, improving overall processing efficiency. The LRU tag is dynamically updated after each access, ensuring that the replacement strategy reflects the actual usage of tag blocks, thus continuously optimizing cache performance over long-term operation.

[0079] In some embodiments, the operation type is a write operation. The command processing unit is specifically used to write the metadata of the scalar command into the target tag block.

[0080] In this embodiment of the invention, if the scalar command's operation type is a write operation, after selecting the target tag block according to the valid flag bit and the Least Recently Used (LRU) policy, the data carried by the scalar command is written to the Data field of the target tag block, and the metadata of the tag block is updated synchronously. This metadata includes: target access address, data type, operation type, wavefront identifier, slot identifier, etc. The Inflight flag bit is set to 0, and the LRU tag position is updated according to the LRU policy.

[0081] When the scalar command operation type is a write operation, no request information to the vector storage module is generated. Therefore, no data is returned by the vector storage module. Only the write operation completion status information is generated as the return information and sent to the wavefront instruction slot module through the e channel so that the corresponding instruction slot in the wavefront instruction slot module is released in order.

[0082] For example, when the scalar command is scmd = write addr 0x40001000 ← 42 and there is no tag block in the data cache module that matches the target access address of the scalar command, the data cache module writes addr = 0x40001000 and data = 42 into the selected target tag block and updates the in-flight flag to 0, indicating that the write operation has been completed.

[0083] In some embodiments, the operation type is a read operation. The command processing unit is specifically used to set the in-flight flag in the target tag block to 1, write the metadata of the scalar command into the target tag block, generate the corresponding request information, and send the request information to the scalar scheduler through the third data transmission channel (channel a).

[0084] In this embodiment of the invention, if the scalar command's operation type is a read operation, after selecting the target tag block according to the valid flag bit and the Least Recently Used (LRU) strategy, the metadata information of the scalar command is written into the metadata field of the target tag block. This metadata includes: target access address, data type, operation type, wavefront identifier, slot identifier, etc., and the LRU flag is updated according to the LRU strategy.

[0085] The data caching module generates a corresponding read request as request information. This request information includes at least the target access address, data type, and read enable flag, and sets the in-flight flag to 1, indicating that the target tag block is waiting for data to return from the vector storage module. The LRU tag position is updated according to the LRU policy. This read request is sent to the scalar scheduler via channel a, which then schedules access to the vector storage module.

[0086] For example, when the scalar command is scmd = read addr 0x40001000 and the data cache module misses, metadata such as addr = 0x40001000 is written into the selected target tag block. At the same time, the corresponding read request is generated and sent, and the Inflight of the tag block is set to 1 to indicate that its data has not yet been returned.

[0087] According to the above embodiments, when the scalar command is a read operation and the data cache module misses, the target tag block is precisely selected based on the valid flag and the LRU strategy. The metadata of the scalar command is written into this target tag block, and a read request containing the necessary access information is generated and sent. This ensures that the read operation request can quickly enter the vector storage module access process while maintaining the effectiveness of the cache replacement strategy. By setting the in-flight flag, the system can clearly identify the return status of the current tag block data, avoiding repeated access or misuse of data that has not been returned, effectively improving data consistency and access efficiency. The above scheme reduces the additional lookup and waiting overhead in the memory access path, shortens the latency from data miss to request issuance, and improves the collaborative efficiency between the data cache module and the vector storage module.

[0088] In some embodiments, the operation type is an atomic operation. Specifically, the command processing unit is used to set the in-flight flag in the target tag block to 1, write the operands in the scalar command into the data field of the target tag block, generate the corresponding request information, and send the request information to the scalar scheduler through the third data transmission channel (channel a).

[0089] In this embodiment of the invention, if the operation type of the scalar command is an atomic operation, after selecting the target tag block according to the valid flag bit and the least recently used strategy, the operands in the scalar command are written into the Data field of the target tag block, and the metadata of the target tag block is updated synchronously. This metadata includes: target access address, data type, operation type, wavefront identifier, and slot identifier.

[0090] Since atomic operations require reading the data corresponding to the target access address from the vector storage module before execution, the data caching module generates a corresponding read request as request information while updating metadata. This request information includes at least the target access address, data type, and read enable flag. Setting the in-flight flag to 1 indicates that the target tag block is waiting for data to return from the vector storage module. This read request is sent to the scalar scheduler via channel A, which schedules access to the vector storage module.

[0091] For example, when the scalar command is scmd = atomic_add addr 0x1000, +5 and the data cache module misses, the operand 5 is written to the Data field of the selected target tag block, along with metadata such as addr = 0x1000. At the same time, a corresponding read request is generated and sent, and the Inflight of the tag block is set to 1 to indicate that the data returned by the vector storage module has not yet been received.

[0092] According to the above embodiment, operands and metadata are written to the target tag block, and a read request for the vector storage module is immediately generated. This allows the data required for atomic operations to be read as quickly as possible even if the cache misses, thereby reducing waiting time. Simultaneously, by setting a flight flag, the status of the target tag block that has not yet returned data is identified, facilitating correct matching and updating after subsequent data returns, ensuring the integrity and data consistency of the atomic operation execution process. This mechanism effectively reduces memory access latency and improves instruction execution efficiency and system concurrency capabilities.

[0093] In some embodiments, the returned information includes first returned information and second returned information. If the operation type is a read operation, the scalar scheduler determines whether the target access address in the request information exists in its write request buffer. If the target access address exists, a data forwarding operation is performed to obtain the first returned information. If the target access address does not exist, the request information is stored in the write request buffer. Based on the storage access cycle permissions, the request information is sent to the vector storage module to obtain the second returned information.

[0094] In this embodiment of the invention, when the scalar scheduler receives a request from the data cache module, it first parses the operation type in the request. If the operation type is a read operation, the scalar scheduler first queries the access address field in the write request buffer. If the write request buffer contains an access address that is the same as the target access address of the read request, it means that the data has been generated by the previous write operation and has not yet been written back to the vector storage module. At this time, the scalar scheduler does not need to send the read request to the vector storage module, but directly reads the corresponding data from the write request buffer, writes it into the return data array as the first return information, and returns it to the data cache module through the d channel to realize the data forwarding operation.

[0095] If there is no record matching the target access address in the write request buffer (i.e., the data cache module misses and the write request buffer misses), the scalar scheduler adds the read request to the write request buffer to wait for scheduling, and sends it to the vector storage module for reading based on the access permission scheduling policy of the vector storage module. The reading result is written as the second return information to the return data array and returned to the data cache module through the d channel.

[0096] For example, suppose the old scalar command is scmd = write addr 0x40001000 ← 42. Because the data caching module performs a cache eviction policy, the data in the tag block corresponding to this write operation is replaced out of the data caching module and stored in the write request buffer as a write request.

[0097] When a new scalar command, `scmd = read addr 0x40001000`, is not found in the data cache module, the scalar scheduler first queries the write request buffer. If the target access address (0x40001000) is hit, the corresponding data value 42 is read directly from the write request buffer and written to the return data array as the return result of the read request. This result is then returned to the data cache module via the `d` channel, thus eliminating the need to access the vector storage module.

[0098] If the write request buffer does not hit the target access address, the scalar scheduler will send the read request to the vector storage module to read the data according to the normal access process. After obtaining the read result, it will write it to the return data array as the return result of the read request, and then return it to the data cache module through the d channel.

[0099] According to the above embodiments, by setting up a write request buffer and a return data array in the scalar scheduler, and combining them with a forwarding mechanism, a fast hit response to read requests is achieved. When the target address accessed by a read request already exists in the write request buffer, the corresponding data can be directly returned from the write request buffer, avoiding access to the vector storage module again, reducing storage access latency and bus occupancy, and improving the system's response speed and bandwidth utilization. Simultaneously, when a miss occurs, read requests are uniformly cached, and access to the vector storage module is scheduled according to access permissions, effectively balancing access conflicts and bandwidth allocation, and improving the overall throughput of the system. Furthermore, in cache replacement scenarios, the above mechanism can maintain the direct accessibility of the replaced data, reducing the overhead of data loss or repeated access, thereby ensuring data consistency and access efficiency while reducing power consumption.

[0100] In some embodiments, the returned information includes second returned information. If the operation type is a write operation, the scalar scheduler stores the request information in the write request buffer. Based on the storage access cycle permissions, the request information is sent to the vector storage module to obtain the second returned information.

[0101] In this embodiment of the invention, if the request information is a write request, and its corresponding operation type is a write operation, the scalar scheduler adds the write request to the write request buffer and waits for scheduling in order. Based on the access permission scheduling policy of the vector storage module, the scalar scheduler sends the write request to the vector storage module to perform the write operation. After the vector storage module completes the write, it writes the write completion confirmation information and related metadata (including the target access address, written data, and completion flag, etc.) as the second return information into the return data array, and returns it to the data cache module sequentially through the d channel.

[0102] According to the above embodiments, write confirmation information can be returned in a timely manner after the write operation is completed, ensuring data consistency and access correctness between the data caching module and the vector storage module, thereby improving the system's reliability and processing efficiency.

[0103] In some embodiments, the vector storage module further includes multiple storage arrays. The scalar access system also includes a vector scheduler (v_sched). When the scalar scheduler has storage access cycle permissions, it further queries the write request buffer based on the target access address of the request information, determines multiple request messages that do not conflict with storage array access, and sends them in parallel to the multiple storage arrays of the vector storage module to obtain a second return message.

[0104] In this embodiment of the invention, the vector storage module is physically divided into multiple independent storage arrays (banks) to support multi-path parallel access. However, due to the physical limitation on the total bandwidth of the vector storage module, there is an upper limit to the number of requests that the vector storage module can receive within a single access scheduling cycle. To achieve a balance between access fairness and throughput between the scalar scheduler and the vector scheduler, the system adopts a weighted round-robin arbitration mechanism during the access arbitration phase, where the arbitration weight of the vector scheduler is slightly higher than that of the scalar scheduler, in order to prioritize the bandwidth allocation requirements for large-scale vector data access.

[0105] The scalar scheduler sequentially scans the write request buffer during each scheduling cycle and determines the target storage array identifier (bank_id) corresponding to the request information based on the target access address and storage array mapping relationship. Internally, the scalar scheduler stores a storage array availability bitmap for the current scheduling cycle. This bitmap includes information on storage arrays already occupied by the vector scheduler, as well as information on storage arrays occupied by request information selected by the scalar scheduler within the current scheduling cycle.

[0106] The scalar scheduler scans from the head of the write request buffer. If a request conflicts with an already occupied storage array in the storage array availability bitmap, it skips the request until it selects the first request with a non-conflicting storage array. The scalar scheduler sends the selected request to the vector scheduler via channel b. It then continues scanning until the bandwidth limit for scalar requests in the current scheduling cycle is reached, or until the vector storage module has no available storage arrays.

[0107] For a read request, after the vector storage module executes the read request and completes the read operation, the read data is used as the second return information and written to the return data array via the C channel. For a write request, after the vector storage module executes the write request and completes the data write operation, the confirmation information and related metadata are used as the second return information and written to the return data array via the C channel.

[0108] The return data array parses the second return information returned by channel c, determining whether the return information corresponds to the request information issued by the scalar scheduler and whether its validity flag (e.g., valid) is in a valid state. If the return data is determined to be valid scalar request data, the data is written to an idle storage slot in the return data array. Following the return order of the scalar requests, the return data array sequentially sends the stored return data to the data cache module via channel d to ensure that the data cache module can update the data of the corresponding tag block and related metadata information in a timely manner.

[0109] According to the above embodiments, by dividing the vector storage module into multiple independent storage arrays and introducing a weighted round-robin arbitration mechanism during the scheduling phase, the scalar scheduler and the vector scheduler can achieve a balance between access fairness and throughput under limited bandwidth conditions.

[0110] In some embodiments, when the scalar scheduler does not have storage access cycle permissions, the scalar scheduler is also used to monitor the occupancy status of each storage array in the vector storage module. If an idle storage array is detected, and the operation type in the request information is the same as the operation type of the storage access operation of the vector scheduler, the request information is inserted into the storage access cycle and sent in parallel to the corresponding storage array to obtain the second return information.

[0111] In this embodiment of the invention, when the scalar scheduler does not obtain exclusive storage access privileges within the current scheduling period, the scalar scheduler monitors in real time the occupancy status of each storage array within the vector storage module. The scalar scheduler determines whether each storage array is currently idle by reading the storage array occupancy bitmap maintained by the vector scheduler. If an idle storage array is detected, and the operation type of the scalar request to be inserted is consistent with the operation type currently being executed by the vector scheduler (i.e., both are read operations or both are write operations), the scalar scheduler can insert the scalar request into the vector scheduling period within that scheduling period to achieve concurrent access scheduling with the vector request.

[0112] After concurrent scheduling is completed, if the scalar request is a read request, its read result data is transmitted to the return data array via channel C as the second return information. If the scalar request is a write request, its write completion confirmation information and related metadata are transmitted to the return data array via channel C as the second return information. The return data array sends the stored return data sequentially to the data cache module via channel D, according to the return order of the scalar requests, to ensure that the data cache module can update the data and related metadata information of the corresponding tag blocks in a timely manner.

[0113] For example, in the Nth scheduling cycle, the vector scheduler issues a vector write request to the vector storage module, occupying banks 0 to 5, while banks 6 and 7 are idle. When the scalar scheduler scans the write request buffer, it detects a scalar write request targeting bank 6, and the operation type of this scalar write request matches the write operation currently being performed by the vector scheduler. Therefore, this scalar write request can be inserted into the vector access flow of the current scheduling cycle for concurrent execution. After the scalar write request is completed, its write confirmation information is stored in the return data array via channel c and finally returned to the data cache module via channel d.

[0114] According to the above embodiments, by dynamically monitoring the storage array occupancy status of the vector scheduler when the scalar scheduler does not obtain exclusive access, and inserting the scalar request into the vector scheduling cycle for concurrent execution when there is a free storage array and the operation type is consistent, efficient utilization of storage bandwidth is achieved, the waiting time of scalar requests is reduced, and storage access latency is effectively reduced. This mechanism not only improves the overall system throughput but also enhances resource sharing and scheduling flexibility between the scalar and vector schedulers, thereby achieving more efficient storage access performance in bandwidth-constrained hardware environments.

[0115] In some embodiments, the wavefront command slot module is further configured to determine the corresponding command slot based on the wavefront identifier and slot identifier in the first operation result and the second operation result, and release the command slot according to a third preset rule.

[0116] In this embodiment of the invention, after all storage access operations of a scalar command are completed, the data caching module generates scalar completion information and transmits it to the wavefront command slot module via the e-channel. The scalar completion information includes a wavefront identifier and a slot identifier, used to identify the position of the scalar command in the wavefront command slot module.

[0117] The wavefront instruction slot module locates the corresponding instruction slot based on the wavefront identifier and slot identifier in the completion information, and performs a slot release operation according to a preset release rule. The release rule is: the current instruction slot can be released when all memory access operations (including read requests, write requests, and atomic operations) contained in the scalar command have been completed, and all instruction slots preceding this slot have been released.

[0118] After the command slot is released, the wavefront command slot module generates the final return information corresponding to the scalar command and sends it to the control queue to notify the control queue that the scalar command has been executed.

[0119] According to the above embodiments, by generating completion information including wavefront and slot identifiers after the scalar command is completed, and combining this with strict slot release rules, ordered submission and precise release control of commands within the wavefront are achieved, effectively avoiding data consistency issues caused by out-of-order release. Introducing the detection of the release status of preceding slots into the release rules ensures that the order of return information received by the control queue is consistent with the command issuance order, thereby improving the predictability of the command submission phase and the stability of system operation.

[0120] In some embodiments, such as Figure 2 As shown, the scalar access system also includes an input queue, a vector output buffer (v_output_buffer), and an atomic operation unit (atomic for short).

[0121] When the input queue receives a vector command (vcmd), it temporarily stores the vector command until the corresponding vector data (v_data) arrives. The vector data includes the target access address and data information.

[0122] Once the vector data arrives, the input queue packages the vector command and the target access address, generates a request message, and sends it to the vector scheduler.

[0123] After receiving a request, the vector scheduler resolves its address. If multiple threads within the same wavefront unit concurrently access the same target address or the same storage array in the vector storage module, an access conflict (bank_conflict) is identified. In this case, the vector scheduler adjusts the execution timing of the threads within the wavefront to avoid access conflicts, thus ensuring the efficiency and consistency of parallel access.

[0124] When the vector storage module receives a request from the vector scheduler, it directly accesses the corresponding storage array and returns the processing result. For write operations, the vector storage module directly writes the data to the corresponding storage array. For read operations, the vector storage module outputs the read vector return data (v_rtn_data) to the vector output buffer (v_output_buffer).

[0125] When the vector scheduler receives an atomic operation request, it forwards it to the atomic operation module. The atomic operation module performs the corresponding atomic update operation or read / write operation at the target access address and returns the operation result to the vector storage module, thereby ensuring the atomicity and consistency of data access.

[0126] After completing data access or atomic operations, the vector storage module generates vector completion information (Vector Done, v_done) and transmits this information to the vector output buffer. The vector output buffer outputs the vector completion information and can also output the read vector return data to indicate the completion of the corresponding vector command execution.

[0127] According to the above embodiments, by establishing an efficient scheduling mechanism among the input queue, vector scheduler, vector storage module, atomic operation module, and vector output buffer, end-to-end optimization of vector commands is achieved. By introducing address conflict detection and wavefront thread scheduling strategies into the vector scheduler, access conflicts caused by concurrent access to the same address or the same storage array can be effectively avoided at the hardware level, thereby improving the parallel efficiency of storage access and the overall system throughput.

[0128] This invention also provides a scalar access method based on a vector memory, as described in the following embodiments. Since the principle behind this method is similar to that of a scalar access system based on a vector memory, its implementation can be found in the implementation of a scalar access system based on a vector memory; repeated details will not be elaborated further.

[0129] This vector memory-based scalar access method, applied to the aforementioned vector memory-based scalar access system, includes:

[0130] The data caching module checks if a tag block exists in the received scalar command's tag block that matches the target access address of the scalar command. If not, it determines the target tag block based on the metadata in the tag block and generates the corresponding request information based on the operation type of the scalar command.

[0131] The scalar scheduler queries the operation type in the request information, accesses its write request buffer or vector storage module based on the operation type, obtains the return information, and sends the return information to the data cache module through the first data transmission channel.

[0132] The data caching module performs a storage access operation on the target tag block corresponding to the returned information to obtain the first operation result. The first operation result is then transmitted to the wavefront command slot module through the second data transmission channel.

[0133] In some embodiments, the returned information includes first returned information and second returned information. If the operation type is a read operation, the scalar scheduler determines whether the target access address in the request information exists in its write request buffer. If the target access address exists, a data forwarding operation is performed to obtain the first returned information. If the target access address does not exist, the request information is stored in the write request buffer. Based on the storage access cycle permissions, the request information is sent to the vector storage module to obtain the second returned information.

[0134] In some embodiments, the returned information includes second returned information. If the operation type is a write operation, the scalar scheduler stores the request information in the write request buffer. Based on the storage access cycle permissions, the request information is sent to the vector storage module to obtain the second returned information.

[0135] In some embodiments, the vector storage module further includes multiple storage arrays. When the scalar scheduler has storage access cycle permissions, it queries the write request buffer based on the target access address of the request information, determines multiple request information that do not have storage array access conflicts, and sends them in parallel to the multiple storage arrays of the vector storage module to obtain the second return information.

[0136] In some embodiments, the system further includes a vector scheduler. When the scalar scheduler does not have storage access cycle permissions, the scalar scheduler monitors the occupancy status of each storage array in the vector storage module. If it detects a storage array that is idle, and the operation type in the request information is the same as the operation type of the storage access operation of the vector scheduler, it inserts the request information into the storage access cycle and sends it in parallel to the corresponding storage array to obtain the second return information.

[0137] In some embodiments, the data caching module includes a tag block determination unit and a command processing unit.

[0138] The tag block determination unit determines the target tag block based on the metadata in the tag block and according to the first preset rule.

[0139] Based on the operation type in the scalar command, the command processing unit performs metadata update operations on the target tag block according to the second preset rule. When the operation type is a read operation or an atomic operation, the corresponding request information is generated and sent to the scalar scheduler.

[0140] In some embodiments, the metadata includes a validity flag. The tag block determination unit queries the validity flag of the tag block. If there is an idle tag block, the tag block with the longest idle time is selected as the target tag block. If there is no idle tag block, the target tag block is selected using the least recently used algorithm, and request information is generated based on the metadata in the target tag block, which is then sent to the scalar scheduler through a third data transmission channel.

[0141] In some embodiments, the operation type is a write operation. The command processing unit writes the metadata of the scalar command into the target tag block.

[0142] In some embodiments, the operation type is a read operation. The command processing unit sets the in-flight flag in the target tag block to 1, writes the metadata of the scalar command into the target tag block, generates request information, and sends the request information to the scalar scheduler through the third data transmission channel.

[0143] In some embodiments, the operation type is an atomic operation. The command processing unit sets the in-flight flag in the target tag block to 1, writes the operands of the scalar command into the data field of the target tag block, generates the corresponding request information, and sends the request information to the scalar scheduler through the third data transmission channel.

[0144] In some embodiments, the wavefront command slot module includes multiple wavefront units, and each wavefront unit includes multiple command slots. The wavefront command slot module receives scalar commands and queries the occupancy status of all command slots in the wavefront unit corresponding to the wavefront identifier based on the wavefront identifier in the scalar command. If there is an idle command slot, the scalar command is stored in that command slot.

[0145] In some embodiments, if a tag block exists that has the same target access address as the scalar command, the data caching module performs a storage access operation on the tag block to obtain a second operation result. The storage access operation includes read operations, write operations, and atomic operations.

[0146] In some embodiments, the wavefront command slot module determines the corresponding command slot based on the wavefront identifier and slot identifier in the first operation result and the second operation result, and releases the command slot according to the third preset rule.

[0147] Figure 3 This is a schematic diagram of the physical structure of an electronic device provided in an embodiment of the present invention, such as... Figure 3 As shown, the electronic device 300 includes a processor 301, a memory 302, and a bus 303.

[0148] The processor 301 and the memory 302 communicate with each other via the bus 303.

[0149] The processor 301 is used to call program instructions in the memory 302 to execute the methods provided in the above-described method embodiments.

[0150] This invention also provides a computer-readable storage medium storing a computer program that, when executed by a processor, implements the above-described scalar access method based on vector memory.

[0151] This invention also provides a computer program product, which includes a computer program that, when executed by a processor, implements the above-described scalar access method based on vector memory.

[0152] Those skilled in the art will understand that embodiments of the present invention can be provided as methods, systems, or computer program products. Therefore, the present invention can take the form of a completely hardware embodiment, a completely software embodiment, or an embodiment combining software and hardware aspects. Furthermore, the present invention can take the form of a computer program product embodied on one or more computer-usable storage media (including, but not limited to, disk storage, CD-ROM, optical storage, etc.) containing computer-usable program code.

[0153] This invention is described with reference to flowchart illustrations and / or block diagrams of methods, apparatus (systems), and computer program products according to embodiments of the invention. 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 illustrations and / or block diagrams. Figure 1 One or more processes and / or boxes Figure 1 A device that provides the functions specified in one or more boxes.

[0154] 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.

[0155] 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.

[0156] The specific embodiments described above further illustrate the purpose, technical solution, and beneficial effects of the present invention. It should be understood that the above descriptions are merely specific embodiments of the present invention and are not intended to limit the scope of protection of the present invention. Any modifications, equivalent substitutions, improvements, etc., 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 scalar access system based on a vector memory, characterized in that, include: Data caching module, wavefront instruction slot module, scalar scheduler and vector storage module; The data caching module is used to query whether there is a tag block in the received scalar command whose target access address is the same as that of the scalar command. If it does not exist, determine the target tag block based on the metadata in the tag block, and generate the corresponding request information based on the operation type of the scalar command; The scalar scheduler includes a write request buffer. The scalar scheduler is used to query the operation type in the request information, access the write request buffer or the vector storage module based on the operation type, obtain return information, and send the return information to the data cache module through the first data transmission channel. The data caching module is used to perform a storage access operation on the target tag block corresponding to the returned information to obtain a first operation result; and transmits the first operation result to the wavefront instruction slot module through the second data transmission channel. The wavefront command slot module is also used to determine the corresponding command slot based on the wavefront identifier and slot identifier in the first operation result and the second operation result, and release the command slot according to the third preset rule.

2. The system according to claim 1, characterized in that, The returned information includes first returned information and second returned information; if the operation type is a read operation, the scalar scheduler determines whether its write request buffer contains the target access address in the request information; if the target access address exists, it performs a data forwarding operation to obtain the first returned information. If the target access address does not exist, the request information is stored in the write request buffer; based on the storage access cycle permission, the request information is sent to the vector storage module to obtain the second return information.

3. The system according to claim 1, characterized in that, The returned information includes second returned information. If the operation type is a write operation, the scalar scheduler stores the request information in the write request buffer; based on the storage access cycle permission, the request information is sent to the vector storage module to obtain the second returned information.

4. The system according to claim 2 or 3, characterized in that, The vector storage module further includes multiple storage arrays; when the scalar scheduler has storage access cycle permissions, the scalar scheduler is also used to query the write request buffer according to the target access address of the request information, determine multiple request information that do not have storage array access conflicts, and send them in parallel to the multiple storage arrays of the vector storage module to obtain the second return information.

5. The system according to claim 2 or 3, characterized in that, Also includes: Vector scheduler; when the scalar scheduler does not have storage access cycle permissions, the scalar scheduler is also used to monitor the occupancy status of each storage array in the vector storage module; If an idle storage array is detected, and the operation type in the request information is the same as the operation type of the storage access operation of the vector scheduler, the request information is inserted into the storage access cycle and sent in parallel to the corresponding storage array to obtain the second return information.

6. The system according to claim 1, characterized in that, The data caching module includes: The tag block determination unit is used to determine the target tag block based on the metadata in the tag block according to a first preset rule; The command processing unit is used to perform metadata update operations on the target tag block according to the operation type in the scalar command and a second preset rule. When the operation type is a read operation or an atomic operation, the unit generates corresponding request information and sends it to the scalar scheduler.

7. The system according to claim 6, characterized in that, The metadata includes: a valid flag bit; the tag block determination unit is specifically used to query the valid flag bit of the tag block, and if there is a tag block in an idle state, select the tag block with the longest idle time as the target tag block; if there is no tag block in an idle state, select the target tag block using the least recently used algorithm, and generate the request information based on the metadata in the target tag block, and send it to the scalar scheduler through the third data transmission channel.

8. The system according to claim 6, characterized in that, The operation type is a write operation; the command processing unit is specifically used to write the metadata of the scalar command into the target tag block.

9. The system according to claim 6, characterized in that, The operation type is a read operation; the command processing unit is specifically used to set the in-flight flag bit in the target tag block to 1, write the meta-information in the scalar command into the target tag block, generate the request information, and send the request information to the scalar scheduler through the third data transmission channel.

10. The system according to claim 6, characterized in that, The operation type is atomic operation; the command processing unit is specifically used to set the in-flight flag bit in the target tag block to 1, write the operand in the scalar command into the data field of the target tag block, generate the corresponding request information, and send the request information to the scalar scheduler through the third data transmission channel.

11. The system according to claim 1, characterized in that, The wavefront command slot module includes multiple wavefront units, and each wavefront unit includes multiple command slots. The wavefront command slot module is used to receive the scalar command, query the occupancy status of all command slots in the wavefront unit corresponding to the wavefront identifier based on the wavefront identifier in the scalar command, and if there is an idle command slot, store the scalar command in that command slot.

12. The system according to claim 1, characterized in that, If a tag block exists that has the same target access address as the scalar command, the data caching module performs a storage access operation on the tag block to obtain a second operation result; wherein the storage access operation includes at least one of the following: a read operation, a write operation, and an atomic operation.

13. A scalar access method based on a vector memory, characterized in that, Applied to the system according to any one of claims 1 to 12, comprising: The data caching module queries the received scalar command to see if there is a tag block in its tag block that has the same target access address as the scalar command; if not, it determines the target tag block based on the metadata in the tag block and generates corresponding request information based on the operation type of the scalar command. The scalar scheduler includes a write request buffer. The scalar scheduler queries the operation type in the request information, accesses the write request buffer or the vector storage module based on the operation type, obtains return information, and sends the return information to the data cache module through the first data transmission channel. The data caching module performs a storage access operation on the target tag block corresponding to the returned information to obtain a first operation result; the first operation result is then transmitted to the wavefront instruction slot module through the second data transmission channel.

14. A computer device comprising a memory, a processor, and a computer program stored in the memory and executable on the processor, characterized in that, When the processor executes the computer program, it implements the method of claim 13.

15. A computer-readable storage medium, characterized in that, The computer-readable storage medium stores a computer program that, when executed by a processor, implements the method of claim 13.

16. A computer program product, characterized in that, The computer program product includes a computer program that, when executed by a processor, implements the method of claim 13.

Images