Data loading method and device, processor, electronic equipment and storage medium

Abstract

The present application relates to the technical field of data processing, and provides a data loading method and device, a processor, electronic equipment and a storage medium, wherein the method comprises: loading a data block from a first memory to a second memory based on a first data loading instruction; before the execution of the first data loading instruction is completed, emitting a second data loading instruction and storing a loading request corresponding to the second data loading instruction in a hardware queue; based on hardware detection logic, detecting the state of a synchronization signal associated with the execution process of the first data loading instruction, and triggering the execution of the loading request stored in the hardware queue in the case where the state of the synchronization signal is detected to have changed. By adopting the mechanism of emitting the second data loading instruction in advance and triggering asynchronously by the hardware queue, the present application realizes fine-grained pipeline overlap of data loading across stages, thereby improving the overall efficiency of data loading and the utilization rate of computing resources.

Description

TECHNICAL FIELD

[0001] The present application relates to the technical field of data processing, and in particular, to a data loading method and device, a processor, an electronic device, and a storage medium. BACKGROUND

[0002] In high-performance computing applications, in order to improve computing performance, a hierarchical storage architecture is usually adopted. When data loading is performed under this architecture, the data needs to be first loaded from a global memory with large capacity but slow speed to a shared memory on a chip with high speed, and then loaded from the shared memory to a register directly accessed by a computing unit.

[0003] This data loading process is usually performed in a display synchronization mode, that is, the processor first transmits a loading instruction to move a data block from the global memory to the shared memory, and then inserts a display synchronization instruction (such as a memory barrier) to wait until the data block is completely written to the shared memory, and then transmits the next loading instruction to load the data block from the shared memory to the register.

[0004] This display synchronization instruction creates a forced synchronization point. For a thread bundle executing the instruction, all subsequent instructions (such as the instruction to load data from the shared memory to the register) dependent on the synchronization event must wait until the entire data block is completely loaded to the shared memory. In this blocking mode, the processor cannot transmit or process the subsequent instruction (such as the instruction to load data from the shared memory to the register) in advance, resulting in an idle cycle in the pipeline and reducing the utilization of computing resources and the overall processing efficiency. SUMMARY

[0005] The present application provides a data loading method, device, processor, electronic device, and storage medium to solve the defect of low computing efficiency caused by the blocking of the data loading pipeline due to the use of display synchronization instructions in the related art and the inability to fully hide memory access delay.

[0006] The present application provides a data loading method, comprising:

[0007] loading a data block from a first memory to a second memory based on a first data loading instruction;

[0008] transmitting a second data loading instruction before the execution of the first data loading instruction is completed, and storing a loading request corresponding to the second data loading instruction in a hardware queue, the second data loading instruction being used to load the data block from the second memory to a register;

[0009] detecting, based on hardware detection logic, a state of a synchronization signal associated with an execution process of the first data load instruction, and triggering execution of the load request stored in the hardware queue in a case where a change in the state of the synchronization signal is detected.

[0010] According to the data load method provided by the present application, the second data load instruction comprises an asynchronous identifier and a phase identifier, the asynchronous identifier is used to represent that the second data load instruction is an asynchronous load instruction, and the phase identifier is used to indicate an expected state of the synchronization signal associated with the second data load instruction.

[0011] According to the data load method provided by the present application, the storing of the load request corresponding to the second data load instruction into the hardware queue comprises:

[0012] The load request, the phase identifier, and a memory barrier address corresponding to the synchronization signal associated with the second data load instruction are stored into one entry of the hardware queue, and a ready flag of the entry is set to an unready state.

[0013] According to the data load method provided by the present application, the synchronization signal is a memory barrier signal, and the method further comprises:

[0014] In a case where the execution of the first data load instruction is completed, a phase of a memory barrier associated with the synchronization signal is flipped to switch the state of the synchronization signal from a first state to a second state.

[0015] According to the data load method provided by the present application, the triggering of the execution of the load request stored in the hardware queue in a case where the change in the state of the synchronization signal is detected comprises:

[0016] In a case where the state of the synchronization signal is switched from the first state to the second state, a ready flag of an entry of the hardware queue associated with the synchronization signal is set to a ready state, and the load request with the ready flag in the ready state is executed.

[0017] According to the data load method provided by the present application, the method further comprises:

[0018] In a case where the execution of the second data load instruction is completed, the load request is removed from the hardware queue, and a phase of a memory barrier associated with the synchronization signal is flipped to switch the state of the synchronization signal from the second state to the first state.

[0019] The present application further provides a data load device, comprising:

[0020] an instruction transmitting unit configured to transmit a first data load instruction and a second data load instruction, wherein the first data load instruction is used to load a data block from a first memory to a second memory, the second data load instruction is used to load the data block from the second memory to a register, and the second data load instruction is transmitted before the first data load instruction is executed completely;

[0021] a hardware queue configured to store a load request corresponding to the second data load instruction;

[0022] a synchronization signal detector configured to detect a state of a synchronization signal associated with an execution process of the first data load instruction, and generate a trigger signal to trigger the load request stored in the hardware queue to be executed if a change of the state of the synchronization signal is detected.

[0023] According to the data load apparatus provided by the present application, the hardware queue and the synchronization signal detector are integrated in a load storage cache unit.

[0024] According to the data load apparatus provided by the present application, the hardware queue comprises a plurality of entries, each of which corresponds to a load request, and each of which comprises a request information field, a ready flag and a barrier address, wherein the request information field is used to store information of the load request, the ready flag is used to indicate whether the corresponding load request can be executed, and the barrier address is used to identify a synchronization signal associated with the load request.

[0025] According to the data load apparatus provided by the present application, the synchronization signal detector is a barrier phase detector, which is configured to:

[0026] continuously detect a phase state of at least one memory barrier;

[0027] if a phase flip of any memory barrier is detected, find an entry in the hardware queue in which the barrier address matches the address of the any memory barrier, and set the ready flag of the found entry to a ready state.

[0028] The present application further provides a processor comprising a first memory and at least one processing unit, each of which comprises a second memory, a register and the data load apparatus as described in any of the above.

[0029] The present application further provides an electronic device comprising a memory, a processor and a computer program stored in the memory and running on the processor, wherein the processor implements the data load method as described in any of the above when executing the computer program.

[0030] The application further provides a non-transitory computer-readable storage medium, which stores a computer program, and the computer program is executed by a processor to implement any one of the data loading methods.

[0031] The application further provides a computer program product, which comprises a computer program, and the computer program is executed by a processor to implement any one of the data loading methods.

[0032] The data loading method, device, processor, electronic device and storage medium provided by the application can transmit the second data loading instruction in advance before the execution of the first data loading instruction is completed, and the second data loading instruction will not be executed immediately or block the processor, but the corresponding loading request of the second data loading instruction is temporarily stored in a dedicated hardware queue, so that the processor can continue to execute subsequent other instructions after the transmission of the instruction, thereby avoiding the pipeline stall caused by waiting for data. At the same time, the hardware automatically detects the synchronization signal associated with the execution process of the first data loading instruction, and when the state of the synchronization signal changes, the loading request temporarily stored in the hardware queue is automatically triggered to load the data from the second memory to the register. The mode of the present application, that is, the instruction transmission in advance, the request for hardware caching and the automatic triggering of the signal, changes the processor software level blocking caused by the display synchronization instruction in the related art into the non-blocking hardware level asynchronous scheduling, not only realizes the fine-grained pipeline overlap of the data loading across stages, but also fully hides the memory access delay, reduces the idle of the computing unit caused by data waiting, improves the instruction transmission throughput, and thus improves the overall efficiency of the data loading and the utilization rate of the computing resource. BRIEF DESCRIPTION OF DRAWINGS

[0033] In order to more clearly illustrate the technical solutions in the application or the related art, the following will briefly introduce the drawings needed to be used in the embodiments or the related art description. Obviously, the drawings in the following description are some embodiments of the application, and other drawings can be obtained by those skilled in the art without any creative effort on the basis of these drawings.

[0034] Figure 1 is a schematic diagram of a data loading pipeline in the related art;

[0035] Figure 2 is a flowchart of the data loading method provided by the application;

[0036] Figure 3 is a schematic diagram of a data loading pipeline provided by the application;

[0037] Figure 4 is a structural schematic diagram of a hardware queue provided by the application;

[0038] Figure 5 is a structural schematic diagram of a data loading device provided by the present application;

[0039] Figure 6 is a structural schematic diagram of a loading storage cache unit provided by the present application;

[0040] Figure 7 is a structural schematic diagram of a processor provided by the present application;

[0041] Figure 8 is a structural schematic diagram of an electronic device provided by the present application. DETAILED DESCRIPTION

[0042] In order to make the objects, technical solutions and advantages of the present application clearer, the technical solutions in the present application will be described clearly and completely below with reference to the drawings in the present application. Obviously, the described embodiments are some embodiments of the present application, rather than all the embodiments. Based on the embodiments in the present application, all other embodiments obtained by those skilled in the art without creative labor fall within the protection scope of the present application.

[0043] In the field of modern high-performance computing, especially in the field of general-purpose graphics processing units (GPGPU), efficient data loading is a key link to determine the overall computing performance. Processors usually adopt a hierarchical storage architecture, which includes a global memory with slow speed but huge capacity, such as high bandwidth memory (HBM), a shared memory with fast speed but limited capacity, such as group shared memory (GSM), and a register with the fastest speed but the smallest capacity, such as thread local register (TLR). The computing unit directly uses the data in the register for operation. Therefore, a typical computing task needs to move the required data block (usually referred to as tile) from the global memory to the shared memory step by step, and then load the data from the shared memory to the register, which constitutes a data loading pipeline.

[0044] In the related art, GPGPU usually adopts an explicit synchronization mode to manage the data loading pipeline. Figure 1 is a schematic diagram of a data loading pipeline in the related art, as shown in Figure 1 , the process usually includes the following steps:

[0045] Step S1, the processor transmits a data loading instruction, such as a Tensor Memory Accel (TMA) instruction, for loading a data block from a global memory (such as HBM) to a shared memory (such as GSM);

[0046] Step S2, an explicit memory barrier instruction, such as mbarrier.wait instruction, is set. The memory barrier is a mandatory synchronization point, and the thread bundle (warp) executing the instruction must pause (stall) execution and wait for the associated memory operation (i.e. TMA instruction) to complete;

[0047] Step S3, when the TMA instruction completely writes the data block to the GSM and notifies the processor through the memory barrier mechanism, the blocked thread bundle is awakened. Subsequently, the processor transmits another data loading instruction (such as ldmatrix instruction) to load the prepared data from the shared memory (such as GSM) to the target register (such as TLR).

[0048] Step S4, after the data is loaded into the register, the computing unit starts executing the computing task, such as matrix multiply accumulate (MMA) calculation. At the same time, the processor can again transmit a TMA instruction to load the next data block from HBM to GSM.

[0049] However, the above data loading process based on the explicit synchronization mode has the following significant defects:

[0050] Firstly, the mbarrier.wait instruction is a mandatory synchronization point, which means that during the entire process of moving data from global memory to shared memory by TMA instruction, all subsequent instructions dependent on the synchronization event must wait for the thread bundle executing the instruction, which will cause bubbles in the pipeline, resulting in idle computing unit and thus significantly reducing the instruction throughput and overall computing efficiency of the processor.

[0051] Secondly, the overlap granularity between data loading and computation is limited. In the above pipeline, the overlap between computation and data loading mainly occurs between the stage of loading a current data block from global memory to shared memory and the computation task of the previous data block, i.e. the TMA instruction execution stage of the current data block overlaps with the MMA stage of the previous data block. However, for a certain data block, the data loading stage (executed by the ldmatrix instruction) from shared memory to register must wait until the stage of loading a complete data block from global memory to shared memory (executed by the TMA instruction) is completed. This means that the ldmatrix instruction must be transmitted after the TMA instruction is executed. This coarse-grained waiting mechanism in units of data blocks limits the deeper pipeline overlap between different data loading stages, and the ability to hide data access delay is insufficient.

[0052] To this end, the present application provides a data loading method and device, which pre-transmits a second data loading instruction during the execution of a first data loading instruction, and caches the request to a hardware queue, and then triggers the execution of the cached request by automatically detecting the state change of a synchronization signal. This realizes the cross-stage overlap between the first-stage data loading and the second-stage data loading, effectively hides the delay of loading data from the second memory to the register, shortens the critical path of data loading, reduces the idle time of the computing unit due to data waiting, and thus improves the overall efficiency of data loading and the utilization rate of computing resources of the processor.

[0053] Figure 2 The present application provides a flowchart of the data loading method, as shown in Figure 2 The method comprises the following steps:

[0054] In step S10, a data block is loaded from a first memory to a second memory based on a first data loading instruction.

[0055] It should be noted that the method provided by the embodiments of the present application can be applied to a data loading device. In a specific computing scenario, for example, when performing MMA operation in GPGPU, the computing unit needs to obtain a large amount of computing data from the external memory. This step is the first stage of the data loading process.

[0056] Specifically, the above-mentioned data block can be understood as a block of data required for a core computing task (for example, an MMA computing iteration). The size of the data block can be configured according to the specific application scenario and hardware capability.

[0057] The first memory generally refers to a storage device with relatively slow speed but large capacity. For example, the first memory can be HBM, which is used as global memory. The second memory refers to a storage device with relatively fast speed but small capacity, which is closer to the core of the computing unit. In the GPGPU architecture, the second memory can be the on-chip shared memory, such as GSM. It is usually managed by software and used to cache the data frequently accessed by the computing unit to reduce the access delay to the first memory. In other architectures, the second memory can also be part of the cache, such as L2 cache.

[0058] The first data loading instruction is an instruction for starting the transmission process of the data block from the first memory to the second memory. For example, in a GPGPU supporting tensor calculation, the instruction can be a TMA instruction. After the instruction is launched, the corresponding data loading operation can be performed to read the data block from the HBM and transfer it to the GSM.

[0059] In step S20, before the execution of the first data loading instruction is completed, a second data loading instruction is launched, and a loading request corresponding to the second data loading instruction is stored in the hardware queue, the second data loading instruction being used to load the data block from the second memory to the register.

[0060] It should be noted that in the conventional data loading process, the second data loading instruction for loading data from the second memory to the register can be launched only after the first data loading instruction (such as TMA instruction) is completely executed, that is, the data block has been completely written into the second memory (such as GSM). This serial execution method introduces unnecessary waiting period, i.e. stall phase, which reduces the efficiency of the pipeline.

[0061] The embodiment of the present application breaks this limitation. After the first data loading instruction is launched, and before the physical data transfer process is completed, the second data loading instruction can be launched in advance. The purpose of the instruction is to further load the same data block loaded into the second memory (such as GSM) into the register (such as TLR) which can be directly accessed by the computing unit.

[0062] Since the first data load instruction is not completed when the second data load instruction is issued, it means that the data block required by the second data load instruction has not arrived completely in the second memory, so the instruction cannot be executed immediately. To solve this problem, the embodiment of the present application uses a hardware queue to temporarily store the load request corresponding to the instruction. The hardware queue is a dedicated hardware buffer structure, which can be an asynchronous load queue (ALQ) in particular, and it is usually integrated in the load store cache (LSC) unit of the processor.

[0063] It can be understood that the load request contains all the information required for a load operation, such as the source address (in the second memory), the target address (in the register), and the data size, etc. By storing the load request in the hardware queue, the second data load instruction has been issued logically, but its physical execution is delayed, so that the explicit waiting (such as through the mbarrier.wait instruction) at the software level is changed into implicit waiting automatically managed by the hardware.

[0064] Step S30, based on the hardware detection logic, detects the state of the synchronization signal associated with the execution process of the first data load instruction, and triggers the execution of the load request stored in the hardware queue when the state of the synchronization signal is detected to change.

[0065] It should be noted that after the load request is cached to the hardware queue, in order to ensure data consistency, the embodiment of the present application provides a mechanism of automatically triggering according to the state of the synchronization signal to inform the hardware queue that the data it is waiting for is ready, so that the load request is executed in time.

[0066] Specifically, the synchronization signal is a hardware signal used to coordinate the execution sequence between different asynchronous operations. In a specific implementation, the synchronization signal can be a memory barrier signal. The memory barrier itself can be associated with a state, such as a phase. When an asynchronous memory operation (such as a data load caused by a TMA instruction) is completed, the hardware will automatically change (for example, flip) the phase of the memory barrier bound to it. Therefore, the change of the memory barrier phase can be used as a sign of the completion of the first data load instruction.

[0067] To achieve automatic detection and triggering, a special hardware detection logic is provided in the system, such as a barrier phase detector (Mbar Phase Detector). The detector continuously monitors the phase state of one or more memory barriers. When it detects that the phase of a certain memory barrier has changed as expected (for example, from state 0 to state 1), it means that the TMA load operation (i.e. the first data load instruction) associated with the barrier has been completed, and the corresponding data block has been safely stored in the second memory.

[0068] When the phase of a certain memory barrier is detected to change, it means that the synchronization signal state corresponding to the memory barrier has changed, at which time the detection logic will notify the hardware queue. After receiving the notification, the hardware queue will find the load request associated with the synchronization signal (i.e. the memory barrier) in the queue and trigger its execution. The specific process of triggering execution can be: the hardware queue marks the load request as ready state and sends it to the subsequent execution pipeline, and the storage controller (such as the GSM controller) reads the data from the second memory (such as the GSM) according to the request information and writes it to the target register.

[0069] The method provided by the embodiment of the present application transmits the second data load instruction in advance before the execution of the first data load instruction is completed, and the second data load instruction will not be executed immediately or block the processor, but the corresponding load request is temporarily stored in a special hardware queue, so that the processor can continue to execute subsequent other instructions after transmitting the instruction, thereby avoiding the pipeline stall caused by waiting for data. At the same time, the hardware automatically detects the synchronization signal associated with the execution process of the first data load instruction, and when the state of the synchronization signal changes, the load request temporarily stored in the hardware queue is automatically triggered to load the data from the second memory to the register. The mode of the present application of transmitting the instruction in advance, requesting the hardware to cache, and automatically triggering the signal converts the blocking of the processor software level caused by the display synchronization instruction in the related art into the non-blocking hardware level asynchronous scheduling, not only realizes the fine-grained pipeline overlap of the data load across stages, but also fully hides the memory access delay, reduces the idle of the computing unit caused by data waiting, improves the instruction transmission throughput, and thus improves the overall efficiency of data loading and the utilization rate of computing resources.

[0070] In order to more clearly illustrate the beneficial effects brought by the embodiment of the present application, an example is described below. Figure 3 is a schematic diagram of the data load pipeline provided by the present application, as Figure 3As shown, in the pipeline, in order to perform the MMA calculation, first, a TMA instruction (i.e. the first data loading instruction) is issued to load the data block from the global memory to the shared memory. After the TMA instruction is issued, an ldmatrix.async instruction (i.e. the second data loading instruction) can be issued immediately and its loading request is cached in the hardware queue without waiting for the TMA instruction to be executed. Meanwhile, the program does not need to insert the waiting or synchronization instruction as shown, and the hardware automatically monitors the synchronization signal associated with the TMA instruction, and when it detects that the TMA instruction is executed, the hardware will immediately trigger the loading request of the ldmatrix.async instruction in the execution request queue. Obviously, from the above description, it can be seen that, compared with the pipeline shown in Figure 3 Figure 1 It can be seen that, compared with the pipeline shown in

[0071] It can be understood that the cross-stage overlap realized by the embodiment of the present application is essentially to overlap and cover the issuing and queuing time of the ldmatrix instruction by the execution time of the TMA instruction, and at the same time, the software waiting gap between the two in the traditional display synchronization scheme is eliminated. When the TMA instruction is completed, the ldmatrix instruction can be executed immediately in the next cycle, realizing seamless connection of the two loading stages serving the same data block.

[0072] Based on any of the above embodiments, the second data loading instruction has a specific format at the instruction encoding level, so that it can be recognized and specially processed by the instruction decoder. Specifically, the second data loading instruction can include an asynchronous identifier and a phase identifier, wherein the asynchronous identifier is used to represent that the second data loading instruction is an asynchronous loading instruction, and the phase identifier is used to indicate the expected state of the associated synchronization signal.

[0073] Specifically, the asynchronous identifier is one or more bits, which is used to explicitly represent the present instruction as an asynchronous loading instruction. In the traditional instruction set architecture, the loading instruction from the shared memory (i.e. the second memory) to the register is usually synchronous, that is, the data must be obtained immediately when the instruction is executed. By setting the asynchronous identifier (such as async), the instruction decoder can identify that this is a loading instruction that does not need to be executed immediately. This allows the hardware to forward the execution request of the instruction to a dedicated hardware queue (such as ALQ) instead of directly sending it to the execution pipeline, thereby avoiding the pipeline stall caused by the data not being ready.

[0074] ​The phase identifier is also one or more bits in the instruction, which can be represented by adding an expect_phase field in the original instruction field. The role of the identifier is to indicate the expected state of the associated synchronization signal that the current asynchronous load instruction is waiting for. As described in the above embodiment, the synchronization signal can be embodied as a memory barrier signal, and its state can be obtained by detecting the phase change of the memory barrier. In the GPGPU architecture, the phase of the memory barrier usually has only two states, such as 0 and 1, and the phase will flip (from 0 to 1 or from 1 to 0) every time the associated asynchronous operation is completed. Therefore, the phase identifier can be encoded as 0 or 1. For example, when the second data load instruction is transmitted, the phase identifier in the second data load instruction can be set to 1 to explicitly inform the hardware at the instruction level that the condition for triggering the execution of the current load request is that the phase of the associated memory barrier becomes 1.

[0075] In the embodiment of the application, by embedding the asynchronous identifier and the phase identifier in the second data load instruction, the management task of the data dependency relationship is transferred from the software layer to the hardware layer. The software developer or compiler only needs to set the two identifiers according to the data flow and dependency relationship when generating the instruction, and the subsequent waiting, detecting and triggering processes can be automatically completed by the hardware. Compared with using blocking synchronization instructions, this method can more finely express the dependency relationship and allow the processor to execute other instructions without dependency relationship during the waiting period, further improving the instruction-level parallelism.

[0076] Based on any of the above embodiments, the hardware queue includes a plurality of entries, each of the entries corresponding to a load request, each of the entries including a request information field, a ready flag and a barrier address, wherein the request information field is used to store information of the load request, the ready flag is used to indicate whether the corresponding load request can be executed, and the barrier address is used to identify a synchronization signal associated with the load request.

[0077] Specifically, Figure 4 is a structural diagram of the hardware queue provided by the application, as Figure 4 As shown, the hardware queue is not a single register, but includes a plurality of entries. This multi-entry design enables the queue to cache a plurality of load requests at the same time, thereby supporting the initiation of asynchronous loads on a plurality of different data blocks in complex computing tasks, greatly improving the parallelism and pipeline depth of data loading.

[0078] Each entry in the queue corresponds to a load request, and each entry includes at least three key fields, i.e., a request info field, a ready flag, and a mbar address. The request info field is used to store information of a specific load request. The information is core data necessary for performing a data load operation, and can mainly include a source address of data (i.e., a starting address of a data block in a second memory, such as a GSM), a target address of data (i.e., an identifier or address of a target register to which data is to be written, such as a TLR), and a data size (i.e., a size of a data block to be loaded this time). When the entry is triggered for execution, the execution unit will complete the actual data transfer according to the information in the field.

[0079] The ready flag is a state control bit, which is used to indicate whether the corresponding load request can be executed. When a load request is stored in an entry because the data it depends on is not ready, the ready flag is set by hardware to an unready state (e.g., 0). The entry will remain in a waiting state and will not be dispatched by the scheduler of the queue. Only when the prerequisite (i.e., the first-stage data load is completed) is met, the flag will be updated to a ready state (e.g., 1). At this time, the entry becomes executable and waits to be scheduled.

[0080] The mbar address field is used to identify a synchronization signal associated with the load request. Preferably, the synchronization signal is specifically a memory barrier signal, and each memory barrier has a unique address or ID in the system. The field is used to store the address or ID of the memory barrier it depends on. Through the field, hardware can establish a clear correspondence between which load request is waiting for which synchronization event. When the phase state of a certain memory barrier changes, the system can use the barrier address as an index to accurately find the entry in the hardware queue waiting for the barrier.

[0081] Through the above structural design, each entry of the hardware queue becomes a task descriptor that records the task content (request info field), execution condition (mbar address), and current state (ready flag). This structured design makes it possible to automatically manage a large number of suspended asynchronous load requests at the hardware level, and provides a basis for implementing an efficient and accurate event-driven execution model.

[0082] Based on any of the above embodiments, in step S20, the load request corresponding to the second data load instruction is stored in the hardware queue, including:

[0083] The memory barrier address corresponding to the synchronization signal associated with the load request, the phase identifier, and the second data load instruction is stored in an entry of the hardware queue, and a ready flag of the entry is set to an unready state.

[0084] In particular, when a processor receives and decodes a second data load instruction containing an asynchronous identifier and a phase identifier, the hardware logic performs the following operations: first, an idle entry is applied from a hardware queue (e.g., an ALQ) that is composed of multiple such entries, each of which is used to store a deferred load request. Then, load information associated with the instruction is filled into the entry, where the load information includes information related to the load request, the phase identifier, and synchronization signal information (e.g., a memory barrier address) associated with the instruction.

[0085] Next, a ready flag of the entry is set to an unready state. Each entry contains a ready flag, which indicates whether the prerequisites of the load request have been met. When the request is just stored in the queue, the first data load instruction on which the request depends has not been completed, and the data has not arrived at the second memory, so the load request does not have execution conditions. At this time, the hardware sets the ready flag of the entry to an unready state (e.g., sets the bit to 0). The entry in the unready state will not be selected and dispatched for execution by the scheduler of the queue.

[0086] When the hardware detection logic (e.g., a barrier phase detector) detects that the associated memory barrier phase changes and matches the expected phase stored in the entry, the detection logic sends a signal to the hardware queue. After receiving the signal, the hardware queue finds one or more corresponding entries and updates their ready flags from the unready state to the ready state (e.g., sets the bit to 1). Once the entry becomes ready, the internal scheduling logic of the queue can dispatch it to the execution unit at an appropriate time to complete the final data load from the second memory to the register.

[0087] The embodiments of the present application construct an efficient hardware management mechanism by explicitly storing the phase identifier, the memory barrier address, and the ready flag in the entry of the hardware queue. The mechanism can accurately record the triggering conditions of each asynchronous load request and track its state in real time, achieving accurate delayed execution of the load operation. This design fixes the complex synchronization logic in hardware, which not only guarantees the correctness of data loading but also maximizes the reduction of software management overhead and unnecessary performance loss.

[0088] According to any of the above embodiments, the synchronization signal is a memory barrier signal. In the field of high-performance computing, especially in GPGPU architecture, a memory barrier is a lightweight hardware synchronization primitive. It does not directly manipulate data itself, but is used to establish a dependency relationship between different asynchronous operations to establish the execution order. Multiple memory barrier objects can be configured in a system, and each object can be associated with different asynchronous instructions (such as memory load or store instructions).

[0089] Accordingly, the method further comprises:

[0090] In the case where the first data load instruction execution is completed, the phase of the memory barrier associated with the synchronization signal is flipped to switch the state of the synchronization signal from the first state to the second state.

[0091] Specifically, this step reveals the trigger source of the synchronization signal state change. When the first data load instruction (for example, a TMA instruction that loads a data block from HBM to GSM) is physically executed and completed, that is, the data block has been completely and safely written to the second memory (such as GSM), the hardware unit responsible for executing the instruction will automatically perform an operation, that is, flip the phase of the memory barrier associated with the instruction.

[0092] Here, the phase is the specific state representation of the memory barrier. A memory barrier usually has at least two states, which are referred to as the first state and the second state in the embodiments of the present application, and can be binary 0 and 1 in specific implementation. For example, initially, the phase of the memory barrier is 0 (i.e., the first state). When the TMA instruction is executed and completed, the hardware will automatically flip its phase to 1 (i.e., the second state). This phase flipping action is a clear signal broadcast to other hardware units in the system, that is, the data transmission from HBM to GSM associated with the memory barrier has been completed.

[0093] In the embodiments of the present application, the phase flip of the memory barrier is used as a flag indicating the completion of the first data load instruction, which utilizes a hardware-native and extremely low-latency communication mechanism. Compared with polling a memory flag bit through software or using an interrupt synchronization mode, this hardware automatic phase flip mechanism is more efficient and can more timely reflect the data ready state, providing a reliable and fast basis for triggering subsequent operations.

[0094] According to any of the above embodiments, in step S30, in the case where the state of the synchronization signal is detected to change, the execution of the load request stored in the hardware queue is triggered, comprising:

[0095] In case the state of the synchronization signal is switched from the first state to the second state, a ready flag of an entry in the hardware queue associated with the synchronization signal is set to a ready state, and a load request of the ready flag being in the ready state is executed.

[0096] In particular, a hardware detection logic (e.g. barrier phase detector) in the system continuously monitors the phase of memory barriers. When it detects that the phase of a target memory barrier (i.e. the state of a synchronization signal) is switched from a first state (e.g. 0) to a second state (e.g. 1), it means that the prerequisite of the second data load instruction associated with the memory barrier has been satisfied.

[0097] Subsequently, the hardware logic queries all entries in the hardware queue. It looks for those entries that satisfy both of the following conditions: one is that the entry is associated with the memory barrier whose phase just flipped; the other is that the expected phase stored in the entry matches the current new state (i.e. the second state) of the memory barrier. For the matched entry, the hardware automatically sets its ready flag from an unready state (e.g. 0) to a ready state (e.g. 1). This state transition indicates that a pending load request now has the execution condition.

[0098] Finally, a scheduling logic of the hardware queue periodically scans the queue. When it finds an entry whose ready flag is in the ready state, it executes the load request stored in the entry. The execution process includes taking the load request out of the queue and dispatching it to the execution pipeline of the load-store buffer unit, which is responsible for actually reading data from the second memory (e.g. GSM) and writing the data into the target register (e.g. TLR).

[0099] Based on any of the above embodiments, the method further comprises:

[0100] In case the second data load instruction is executed, the load request is removed from the hardware queue, and a phase flip of the memory barrier associated with the synchronization signal is performed to switch the state of the synchronization signal from the second state to the first state.

[0101] In particular, when the second data load instruction is executed, it means that the load-store buffer unit has successfully loaded the data block from the second memory (e.g. GSM) into the target register (e.g. TLR). At this moment, the task of the load request stored in the hardware queue is completed. Therefore, the hardware removes it from the hardware queue. In implementation, it usually means that the entry where the load request is located is marked as free or invalid, so that it can be reused by a subsequent new asynchronous load instruction. This ensures that the limited hardware queue resource can be efficiently recycled.

[0102] Secondly, the hardware flips the phase of the memory barrier associated with the synchronization signal. The phase here is opposite to the phase flipping direction after the execution of the first data loading instruction, i.e., the state of the synchronization signal is switched from the second state (e.g., 1) back to the first state (e.g., 0). The purpose of this reverse flipping is to reset the memory barrier to the initial state, so as to prepare for the next round of data loading and synchronization. For example, in the iteration loop of the matrix multiplication calculation, the data loading is performed in batches and in a pipeline manner. When the loading process (from HBM to register) of the first batch of data is completed, the phase of the memory barrier is flipped back to 0, thereby creating conditions for the loading process of the second batch of data. After the execution of the TMA loading instruction of the second batch of data, the phase of the barrier can be flipped to 1 again, thereby triggering the GSM to register loading of the second batch of data, and so on.

[0103] The embodiment of the present application constructs a complete, recyclable, and pipelined asynchronous data loading closed loop, which not only optimizes the delay of single data loading, but also ensures the continuous and efficient supply of data flow in continuous computing tasks.

[0104] The data loading device provided by the present application is described below. The data loading device described below can be correspondingly referred to the data loading method described above.

[0105] Based on any of the above embodiments, the embodiment of the present application further provides a data loading device, which can implement the data loading method described in any of the above embodiments. Figure 5 is a structural schematic diagram of the data loading device provided by the present application, as Figure 5 shown, the device comprises:

[0106] The instruction transmitting unit 510 is configured to transmit a first data loading instruction and a second data loading instruction, wherein the first data loading instruction is used to load a data block from a first memory to a second memory, the second data loading instruction is used to load the data block from the second memory to a register, and the second data loading instruction is transmitted before the execution of the first data loading instruction is completed;

[0107] The hardware queue 520 is configured to store a loading request corresponding to the second data loading instruction;

[0108] The synchronization signal detector 530 is configured to detect the state of a synchronization signal associated with the execution process of the first data loading instruction, and generate a trigger signal to trigger the execution of the loading request stored in the hardware queue when it is detected that the state of the synchronization signal changes.

[0109] Specifically, the apparatus provided by the embodiments of the present application can include an instruction issue unit, a hardware queue, a synchronization signal detector, and the like. In the front-end pipeline of the processor, the instruction issue unit is responsible for obtaining instructions from the instruction cache, decoding, and issuing them to the corresponding execution units or functional units. Specifically, the instruction issue unit can be used to issue a first data load instruction, such as a TMA instruction, which is used to initiate an asynchronous data transfer to load a data block from a first memory (such as an HBM) with larger capacity and slower speed to a second memory (such as a GSM) closer to the calculation unit and with faster speed.

[0110] In addition, the instruction issue unit can also be used to issue a second data load instruction, such as an ldmatrix.async instruction, which is used to load the same data block loaded into the second memory into a register (such as a TLR) that can be directly accessed by the calculation unit. A key feature is that the instruction issue unit has the ability of pipeline optimization, which can issue the second data load instruction in advance after the first data load instruction is issued and before the physical execution process of the first data load instruction is completed.

[0111] The hardware queue is configured to store a load request corresponding to the second data load instruction. The hardware queue is a dedicated hardware buffer structure, such as an asynchronous load queue (ALQ). When the instruction issue unit issues the second data load instruction in advance, the instruction cannot be executed immediately because the required data is not ready. At this time, the decoded load request (containing source address, target address, data size, and the like) will not be sent to the execution pipeline, but will be stored in the hardware queue for temporary storage and suspension. The existence of the hardware queue is a key physical basis for decoupling instruction issue and execution and hiding data dependency delay.

[0112] The synchronization signal detector is configured to detect the state of a synchronization signal associated with the execution process of the first data load instruction. The detector is a special hardware logic unit, such as a memory barrier phase detector. It continuously monitors the state of one or more synchronization signals (such as the phase of one or more memory barriers).

[0113] Further, the synchronization signal detector is further configured to generate a trigger signal to trigger the execution of the load request stored in the hardware queue when the state of the synchronization signal changes. When the first data load instruction is executed and causes the state of the associated synchronization signal to change as expected (for example, the phase of the memory barrier is reversed), the synchronization signal detector can capture the change and immediately generate a trigger signal. The trigger signal is sent to the hardware queue, notifying that a suspended load request can now be executed. After receiving the signal, the hardware queue will dispatch the corresponding load request to the execution pipeline.

[0114] In the embodiment of the present application, the instruction transmitting unit, the hardware queue and the synchronization signal detector work together to form an efficient hardware automated data loading flow. The instruction transmitting unit is responsible for transmitting data loading instructions in advance, the hardware queue is responsible for temporarily storing loading requests, and the synchronization signal detector plays the role of event monitoring and dispatcher. This structure changes the synchronization waiting that is traditionally completed by software through blocking instructions into an event-driven mechanism that is handled by hardware in parallel, thereby realizing deep overlap of loading operations across storage levels, effectively hiding data loading delay and improving the computing throughput of the entire device.

[0115] Based on any of the above embodiments, the hardware queue and the synchronization signal detector are integrated in a loading storage cache unit.

[0116] Specifically, the loading storage cache (LSC) unit is a crucial component in a processor core, which uniformly manages access requests to on-chip memories at various levels (such as L1 cache, GSM, etc.). All loading, storage and atomic operation instructions from the computing unit are sent to the LSC unit for processing.

[0117] Since the hardware queue (i.e., the asynchronous loading queue) itself is used to manage a special kind of loading request (asynchronous loading from the second memory to the register), placing it in the LSC unit that uniformly manages all loading requests can reuse the existing address calculation, port arbitration and data path hardware resources of the LSC unit, avoiding repeated design and saving chip area and power consumption.

[0118] Similarly, integrating the synchronization signal detector in the LSC unit also has significant advantages. The LSC unit itself is responsible for processing memory barrier instructions, ensuring memory consistency and executing order. Therefore, integrating the synchronization signal detector in the LSC unit allows the detection logic to access the state information of the memory barrier with the lowest delay. More importantly, this close integration allows the synchronization signal detector to immediately and directly send a trigger signal to the asynchronous loading queue that is also located in the LSC unit after detecting a state change. This internal signal transmission path is extremely short and fast, far superior to communication between two independent and physically separated units.

[0119] Figure 6 is a structural schematic diagram of the loading storage cache unit provided by the present application, as Figure 6As shown, the load-store buffer unit mainly consists of three parts: a front-end interface, a buffer control module, and a back-end module. The front-end interface is responsible for receiving instructions sent from the instruction issue unit and performing preliminary analysis and distribution. The buffer control module includes multiple sub-modules, such as a memory barrier controller, a barrier phase detector, an asynchronous load queue (ALQ), a shared memory controller, a barrier state memory, and shared memory, which work together to complete implicit synchronization and data loading.

[0120] Specifically, the asynchronous load queue is mainly used to temporarily store the load requests of the ldmatrix.async instruction that is issued in advance by the instruction issue unit but whose dependent data is not ready. These requests wait in the queue until the data is ready.

[0121] The memory barrier controller is responsible for managing the state of the memory barrier, and it has a barrier phase detector integrated inside. The detector's responsibility is to continuously monitor the barrier phase stored in the memory barrier SRAM (Memory Barrier SRAM). Once it detects that the phase of a certain memory barrier has flipped, it will immediately notify the ALQ. The barrier state memory refers to a special SRAM (Static Random Access Memory) used to store the current state (or phase) of the memory barrier. The barrier phase detector determines the phase change by reading the value here.

[0122] When the load request in the ALQ is activated (i.e., marked as ready) by the barrier phase detector, the ALQ sends the request to the shared memory controller (GSM Controller), which is then responsible for generating the specific address and reading the data from the shared memory. The back-end module is responsible for organizing and formatting the data read from the shared memory and then sending it to the target register through the data bus.

[0123] Based on any of the above embodiments, the synchronization signal detector is a barrier phase detector, which is configured to:

[0124] continuously detect the phase state of at least one memory barrier;

[0125] in the case of detecting the phase flip of any memory barrier, find an entry in the hardware queue that matches the barrier address and the address of the any memory barrier, and set the ready flag of the found entry to the ready state.

[0126] Specifically, the synchronization signal detector can be a barrier phase detector. The detector is a hardware logic specially designed for handling memory barrier based synchronization mechanism. The barrier phase detector is configured to perform the following series of automated operations:

[0127] First, the phase state of at least one memory barrier is continuously detected. Instead of being passively queried when needed, the barrier phase detector functions as a monitoring unit that monitors the phase state memory of one or more memory barriers in real time and in parallel. This continuous hardware monitoring ensures that any phase change can be captured with minimal delay.

[0128] Second, in the event of a phase flip of any memory barrier, the barrier phase detector will take immediate action. The phase flip is a hardware signal indicating the completion of the first stage data loading. Once this signal is captured, the detector proceeds to the next step.

[0129] Next, the barrier phase detector searches the hardware queue for an entry whose barrier address matches that of the memory barrier that just experienced a phase flip. It uses the address of the memory barrier that just experienced a phase flip as an index to perform a quick matching search of all entries in the hardware queue.

[0130] Finally, for the matching entry found, the barrier phase detector (or the logic it controls) sets the ready flag of the found entry to a ready state. That is, it updates the ready flag in these entries from unready (e.g., 0) to ready (e.g., 1).

[0131] The barrier phase detector in the embodiments of the present application can automatically locate all subsequent loading tasks that depend on the data after the completion of the first stage data loading operation and activate them. This mechanism is the key to achieving ultra-low latency data loading, which reduces synchronization overhead from milliseconds or microseconds at the software level to nanoseconds at the hardware level, thereby greatly improving data loading efficiency and overall performance.

[0132] Based on any of the above embodiments, the embodiments of the present application also provide a processor that integrates the aforementioned data loading device at the system architecture level, thereby embedding the data loading method advantages in the hardware and providing efficient data loading capabilities for various compute-intensive applications. Figure 7 is a structural schematic diagram of the processor provided by the present application, as Figure 7 shown, the processor 700 includes a first memory 710 and at least one processing unit 720, each processing unit 720 including a second memory 721, a register 722, and a data loading device 723 as described in any of the above embodiments.

[0133] Specifically, in the embodiments of the present application, the processor can be various types of processing chips for high-performance computing, such as a graphics processing unit (GPU), a general-purpose graphics processing unit (GPGPU), a tensor processing unit (TPU), or an application specific integrated circuit (ASIC) customized for a specific field (such as artificial intelligence), etc.

[0134] The first memory serves as the global memory of the processor and provides large-capacity data sharing and persistent storage for all processing units in the system. In modern processor designs that pursue high throughput, the first memory can be a high-bandwidth memory (HBM).

[0135] The core computing power of the processor is provided by one or more processing units. In the architecture of a GPGPU, a processing unit can be a streaming multiprocessor. A high-performance processor can integrate tens or even hundreds of such processing units to support large-scale parallel computing.

[0136] Further, the interior of each processing unit includes a second memory, a register, and a data loading device. The second memory is an on-chip high-speed memory located in the interior of each processing unit. For example, in a streaming multiprocessor, the second memory can be a shared memory such as a GSM. It serves as a controllable first-level data cache for temporarily storing data frequently accessed or shared by threads in the processing unit, so as to avoid frequent and high-latency access to the external first memory. The register is the fastest storage level in the interior of the processing unit and is directly connected to the computing engine such as an arithmetic logic unit, a floating point unit, or a tensor computing core.

[0137] The interior of each processing unit integrates a data loading device as described in the foregoing embodiments. The device contains all functional logic such as instruction emission, hardware queue management, and synchronization signal detection, and is preferably integrated in the load-store cache unit of the processing unit.

[0138] Under the above processor architecture, a complete data loading process is as follows: when a computing task running on a certain processing unit needs data, the data loading device inside the processing unit coordinates the entire process. It first transmits a first data loading instruction to start the transmission of data from a first memory (such as HBM) at the processor level to a second memory (such as GSM) at the processing unit level. In the transmission process, it transmits a second data loading instruction in advance and caches the corresponding loading request in its internal hardware queue. When the first-stage transmission is completed, the hardware signals through the phase flip of the memory barrier, and the data loading device can detect the signal at the speed of pure hardware and immediately trigger the execution of the cached request in the hardware queue to efficiently load data from the second memory (such as GSM) to the register (such as TLR) for the computing core to use.

[0139] It should be noted that the processor in the embodiment of the application can include other structures in addition to the above structure, and the embodiment of the application does not make specific limitations.

[0140] The processor provided by the embodiment of the application completely hardware and automates the complex asynchronous data loading synchronization mechanism across the storage hierarchy by integrating the data loading device inside each processing unit. This enables the processor to greatly hide data access delay when performing data-intensive tasks, ensures that the computing core can obtain continuous and stable data flow supply, and avoids performance bottlenecks caused by waiting for data.

[0141] Figure 8 An example of a schematic diagram of the physical structure of an electronic device is shown in Figure 8 As shown, the electronic device can include a processor 810, a communication interface 820, a memory 830, and a communication bus 840, wherein the processor 810, the communication interface 820, and the memory 830 complete mutual communication through the communication bus 840. The processor 810 can invoke a logical instruction in the memory 830 to execute a data loading method, which includes: based on a first data loading instruction, loading a data block from a first memory to a second memory; before the execution of the first data loading instruction is completed, transmitting a second data loading instruction, and storing a loading request corresponding to the second data loading instruction into a hardware queue, the second data loading instruction being used to load the data block from the second memory to a register; based on hardware detection logic, detecting the state of a synchronization signal associated with the execution process of the first data loading instruction, and in the case where the state of the synchronization signal changes, triggering the execution of the loading request stored in the hardware queue.

[0142] Further, the logic instructions in the memory 830 described above can be implemented in the form of software functional units and sold or used as independent products, and can be stored in a computer readable storage medium. Based on such understanding, the technical solutions of the present application or the parts that essentially contribute to the related art or the parts of the technical solutions can be embodied in the form of a software product. The computer software product is stored in a storage medium and includes a plurality of instructions for causing a computer device (which can be a personal computer, a server, or a network device, etc.) to execute all or part of the steps of the methods described in the various embodiments of the present application. The aforementioned storage medium includes: a U disk, a mobile hard disk, a read-only memory (ROM, Read-Only Memory), a random access memory (RAM, Random Access Memory), a magnetic disk or an optical disk, and various media that can store program codes.

[0143] In another aspect, the present application also provides a computer program product, which comprises a computer program, the computer program can be stored on a non-transitory computer readable storage medium, and the computer program can be executed by a processor to enable a computer to execute the data loading method provided by the above-mentioned methods. The method comprises: loading a data block from a first memory to a second memory based on a first data loading instruction; before the execution of the first data loading instruction is completed, a second data loading instruction is transmitted, and a loading request corresponding to the second data loading instruction is stored in a hardware queue, the second data loading instruction is used to load the data block from the second memory to a register; based on hardware detection logic, detecting the state of a synchronization signal associated with the execution process of the first data loading instruction, and triggering the execution of the loading request stored in the hardware queue in the case where it is detected that the state of the synchronization signal changes.

[0144] In another aspect, the present application also provides a non-transitory computer readable storage medium, which stores a computer program, and the computer program is executed by a processor to implement the data loading method provided by the above-mentioned methods. The method comprises: loading a data block from a first memory to a second memory based on a first data loading instruction; before the execution of the first data loading instruction is completed, a second data loading instruction is transmitted, and a loading request corresponding to the second data loading instruction is stored in a hardware queue, the second data loading instruction is used to load the data block from the second memory to a register; based on hardware detection logic, detecting the state of a synchronization signal associated with the execution process of the first data loading instruction, and triggering the execution of the loading request stored in the hardware queue in the case where it is detected that the state of the synchronization signal changes.

[0145] The device embodiments described above are merely illustrative, wherein the units described as separate components can or can not be physically separate, and the components displayed as units can or can not be physical units, i.e., can be located in one place, or can be distributed to multiple network units. Part or all of the modules can be selected to achieve the purposes of the embodiments according to actual needs. Those skilled in the art can understand and implement without creative labor.

[0146] Through the description of the above embodiments, those skilled in the art can clearly understand that the embodiments can be realized by means of software and the necessary general hardware platform, and of course can also be realized by hardware. Based on such understanding, the above technical solutions can be embodied in the form of a software product, which can be stored in a computer readable storage medium, such as a ROM / RAM, a magnetic disk, an optical disk, etc., and includes a number of instructions to make a computer device (which can be a personal computer, a server, or a network device, etc.) execute the methods described in each embodiment or some parts of the embodiments.

[0147] Finally, it should be noted that: the above embodiments are only used to illustrate the technical solutions of the present application, and not to limit them; although the present application has been described in detail with reference to the foregoing embodiments, those skilled in the art should understand that: it can still modify the technical solutions recorded in the foregoing embodiments, or make equivalent replacement for part of the technical features; and these modifications or replacements do not make the essence of the corresponding technical solutions deviate from the spirit and scope of the technical solutions of the embodiments of the present application.

Claims

1. A data loading method, characterized in that, include: Based on the first data load instruction, the data block is loaded from the first memory to the second memory; Before the first data loading instruction is completed, a second data loading instruction is issued, and the loading request corresponding to the second data loading instruction is stored in the hardware queue. The second data loading instruction is used to load the data block from the second memory into the register. Based on hardware detection logic, the state of the synchronization signal associated with the execution process of the first data loading instruction is detected, and if the state of the synchronization signal changes, the loading request stored in the hardware queue is triggered to be executed.

2. The data loading method according to claim 1, characterized in that, The second data loading instruction includes an asynchronous identifier and a phase identifier. The asynchronous identifier is used to characterize the second data loading instruction as an asynchronous loading instruction, and the phase identifier is used to indicate the desired state of the second data loading instruction for the associated synchronization signal.

3. The data loading method according to claim 2, characterized in that, The step of storing the loading request corresponding to the second data loading instruction into the hardware queue includes: The loading request, the phase identifier, and the memory barrier address corresponding to the synchronization signal associated with the second data loading instruction are stored in an entry of the hardware queue, and the ready flag of the entry is set to a not ready state.

4. The data loading method according to claim 1, characterized in that, The synchronization signal is a memory barrier signal, and the method further includes: Upon completion of the first data loading instruction, the memory barrier associated with the synchronization signal is phase-flipped to switch the state of the synchronization signal from the first state to the second state.

5. The data loading method according to claim 4, characterized in that, The step of triggering the execution of the loading request stored in the hardware queue when a change in the state of the synchronization signal is detected includes: If the state of the synchronization signal is detected to switch from the first state to the second state, the ready flag of the entry associated with the synchronization signal in the hardware queue is set to the ready state, and a load request with the ready flag set to the ready state is executed.

6. The data loading method according to any one of claims 1 to 5, characterized in that, Also includes: Upon completion of the second data loading instruction, the loading request is removed from the hardware queue, and the memory barrier associated with the synchronization signal is phase-flipped to switch the state of the synchronization signal from the second state to the first state.

7. A data loading device, characterized in that, include: The instruction issuing unit is configured to issue a first data load instruction and a second data load instruction, wherein the first data load instruction is used to load a data block from a first memory to a second memory, the second data load instruction is used to load the data block from the second memory to a register, and the second data load instruction is issued before the first data load instruction is executed. A hardware queue is configured to store load requests corresponding to the second data load instruction; A synchronization signal detector is configured to detect the state of a synchronization signal associated with the execution of the first data loading instruction, and to generate a trigger signal to trigger the execution of the loading request stored in the hardware queue if a change in the state of the synchronization signal is detected.

8. The data loading device according to claim 7, characterized in that, The hardware queue and the synchronization signal detector are integrated in the load storage cache unit.

9. The data loading device according to claim 7, characterized in that, The hardware queue includes multiple entries, each entry corresponding to a loading request. Each entry includes a request information field, a ready flag, and a barrier address. The request information field stores information about the loading request, the ready flag indicates whether the corresponding loading request can be executed, and the barrier address identifies the synchronization signal associated with the loading request.

10. The data loading device according to claim 9, characterized in that, The synchronization signal detector is a barrier phase detector, and the barrier phase detector is configured as follows: Continuously monitor the phase state of at least one memory barrier; If a phase flip of any memory barrier is detected, the hardware queue is searched for an entry whose barrier address matches the address of any memory barrier, and the ready flag of the found entry is set to the ready state.

11. A processor, characterized in that, It includes a first memory and at least one processing unit, each processing unit including a second memory, a register, and a data loading device as described in any one of claims 7 to 10.

12. An electronic device comprising a memory, a processor, and a computer program stored in the memory and running on the processor, characterized in that, When the processor executes the computer program, it implements the data loading method as described in any one of claims 1 to 6.

13. A non-transitory computer-readable storage medium having a computer program stored thereon, characterized in that, When the computer program is executed by a processor, it implements the data loading method as described in any one of claims 1 to 6.

Images