Memory access optimization method and apparatus for collective communication operators and computing device

Abstract

The application provides a memory access optimization method and device for collective communication operators and a computing device, and relates to the technical field of graphic processor computing. The method comprises the following steps: identifying a memory access instruction sequence in a collective communication operator to be optimized, wherein the memory access instruction sequence comprises a standard load instruction and a standard storage instruction; replacing the standard load instruction with a batch load instruction and replacing the standard storage instruction with a batch storage instruction; and configuring the batch load instruction and the batch storage instruction to perform batch operations on continuous memory address spaces at a thread bundle level, and finally executing the batch load instruction and the batch storage instruction, so as to reduce the number of memory access instruction transmissions, reduce the scheduling burden of a thread bundle scheduler, and effectively improve the execution efficiency and bandwidth utilization of the collective communication operator in the distributed training of the graphic processor.

Description

TECHNICAL FIELD

[0001] The present application relates to the technical field of graphics processor computing, and particularly relates to a memory access optimization method and device for collective communication operators and a computing device. BACKGROUND

[0002] In the field of high-performance computing of graphics processing units (GPUs), collective communication operations are the core components of distributed deep learning training. Operators such as AllReduce and AllGather in collective communication libraries are widely used to synchronize model parameters and gradients among multiple GPUs, and the performance of the operators directly determines the scalability and efficiency of large-scale deep learning model training.

[0003] At present, mainstream GPU collective communication libraries mainly rely on traditional load and store instructions for memory access during execution. Each thread is usually responsible for accessing an independent data element. When a large number of threads concurrently access discontinuous memory addresses, non-merged memory access occurs, increasing the number of memory subsystem transactions. In addition, this fine-grained instruction mode requires the emission of a large number of single operation instructions, which places a heavy scheduling burden on the thread bundle scheduler and occupies integer operation pipeline resources. In particular, when dealing with the large block continuous memory access scenario specific to collective communication operators, the efficiency of GPU memory access and instruction execution is low.

[0004] Therefore, there is an urgent need for an optimization scheme that can efficiently and batch access continuous memory blocks to improve the execution efficiency and bandwidth utilization of GPU distributed training. SUMMARY

[0005] The present application provides a memory access optimization method and device for collective communication operators and a computing device to solve the problems of low memory access efficiency, large number of instruction emissions, and heavy burden on the thread bundle scheduler of traditional memory instructions for collective communication operators, which limits the overall performance and bandwidth utilization of GPU distributed training.

[0006] To solve the above technical problems, the present application provides the following technical solutions:

[0007] The application provides a memory access optimization method for collective communication operators, which comprises the following steps: firstly, identifying a memory access instruction sequence contained in a collective communication operator to be optimized, the instruction sequence being composed of standard load instructions and standard store instructions; then, replacing the identified standard load instructions with batch load instructions capable of batch loading data from a continuous memory address space at a thread bundle level; and simultaneously, replacing the standard store instructions with batch store instructions capable of batch storing data to a continuous memory address space at a thread bundle level; by executing the batch load instructions and the batch store instructions, the number of memory access instruction transmissions can be effectively reduced, and the scheduling burden of a thread bundle scheduler can be significantly lightened.

[0008] According to the application, the memory access optimization method for collective communication operators is selected from one or more of an all-reduce operator (AllReduce), a reduce-scatter operator (ReduceScatter), an all-gather operator (AllGather), and an all-to-all operator (AllToAll).

[0009] According to the application, the batch load instructions and the batch store instructions are explicit batch operation instructions directly supported by hardware.

[0010] According to the application, the batch load instructions and the batch store instructions are called through a built-in function of a compiler.

[0011] According to the application, when the built-in function is called, parameters related to a cache strategy need to be configured, the parameters including first-level cache control attributes and second-level cache control attributes.

[0012] According to the application, when the batch store instructions are executed, data is directly stored to a target memory address, and the whole storage process does not need to access any temporary buffer.

[0013] According to the application, when the batch load instructions are executed, a memory access address is calculated through a scalar arithmetic logic unit.

[0014] According to the application, the batch store instructions support a burst transmission mode.

[0015] According to the application, the memory access optimization method for collective communication operators further comprises loop unrolling processing on a code loop containing the batch load instructions and the batch store instructions.

[0016] According to the memory access optimization method for collective communication operators provided by the application, the loop unrolling processing includes backing up the scalar address in the loop control by using the thread bundle state register.

[0017] The application provides a memory access optimization device for collective communication operators, which comprises an identification module, a load instruction replacement module, a storage instruction replacement module and an execution module.

[0018] The application provides a computing device, which comprises one or more graphic processors and a memory for storing one or more programs.

[0019] The application provides a computer readable storage medium, which stores a computer program.

[0020] The application further provides a computer program product comprising a computer program.

[0021] The application provides a memory access optimization method, device and computing equipment for a collective communication operator. BRIEF DESCRIPTION OF DRAWINGS

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

[0023] Figure 1 is a flowchart of a memory access optimization method for a collective communication operator provided by the present application.

[0024] Figure 2 is a performance comparison diagram of a batch memory operation instruction based on BR200 hardware provided by the present application.

[0025] Figure 3 is a structural diagram of a memory access optimization device for a collective communication operator provided by the present application.

[0026] Figure 4 is a structural diagram of a computing equipment provided by the present application. DETAILED DESCRIPTION

[0027] In order to make the purpose, technical solutions and advantages of the present application clearer, the technical solutions in the present application will be described clearly and completely in combination with the drawings in the present application. Obviously, the described embodiments are some embodiments of the present application, not all embodiments. Based on the embodiments in the present application, all other embodiments obtained by those skilled in the art without any creative effort belong to the protection scope of the present application.

[0028] In the technical field to which the present application pertains, GPUs usually employ single instruction multiple threads (SIMT) architecture to perform computing tasks. Under this architecture, multiple threads are organized into thread warps as the basic scheduling and execution unit. Threads in each thread warp execute the same instruction synchronously, but can process different data.

[0029] In terms of memory subsystem, GPUs usually contain multiple levels of cache; at the programming model level, GPU programming languages such as compute unified device architecture (CUDA) provide rich memory access instructions. Traditional load and store instructions usually operate at thread granularity, and each thread accesses its corresponding data element through independent instructions. The compiler optimizes these memory access instructions during compilation, including but not limited to loop unrolling and other conventional optimization methods, to reduce instruction overhead and improve instruction-level parallelism.

[0030] To optimize memory access performance, the present application provides a memory access optimization method, device and computing equipment for collective communication operators. By merging memory access mechanisms, when multiple threads in a thread warp access consecutive memory addresses, these access requests are merged into fewer memory transactions, thereby improving memory bandwidth utilization.

[0031] Figure 1 is a flowchart of a memory access optimization method for collective communication operators provided by the present application.

[0032] As shown in Figure 1 , the method comprises the following steps:

[0033] S101, identifying a memory access instruction sequence in a collective communication operator to be optimized.

[0034] In an embodiment of the present application, the memory access instruction sequence comprises standard load instructions and standard store instructions.

[0035] By way of example, the standard load instruction refers to a traditional memory read instruction used in GPU programming. For example, the ld.global instruction in the CUDA programming model. The characteristic of this type of instruction is that each thread independently performs a load operation, and each thread usually accesses only one data element.

[0036] Specifically, when multiple threads access non-contiguous memory addresses, it will cause non-coalesced access, increasing the number of memory requests. For example, in the gradient summation process of implementing the AllReduce operator, when each GPU thread reads the local gradient value using the standard load instruction, multiple scattered memory transaction requests will be generated due to the non-contiguous access addresses between threads.

[0037] By way of example, the standard store instruction refers to the traditional memory write instruction used in GPU programming. For example, the st.global instruction in the CUDA programming model. Similar to the standard load instruction, each thread independently performs the store operation to write data to the respective target address.

[0038] Specifically, this storage method will generate a large number of storage transactions when multiple threads simultaneously write to non-contiguous memory regions, reducing the utilization of storage bandwidth. For example, after completing the gradient reduction calculation, each thread uses the standard store instruction to write the result to the global memory, and due to the dispersion of the storage addresses, the burst transfer capability of the memory controller cannot be fully utilized.

[0039] In an embodiment of the present application, the collective communication operator includes at least one of an AllReduce operator, a ReduceScatter operator, an AllGather operator, and an AlltoAll operator.

[0040] Specifically, the AllReduce operator is an operator used in distributed training to perform reduction operations (such as summation, averaging, etc.) on tensor data on all computing nodes and synchronize the final result to all nodes. For example, in deep learning training, the AllReduce operator is used to aggregate model gradients on multiple GPU cards to ensure consistency of model parameters.

[0041] Specifically, the ReduceScatter operator is an operator that first performs reduction operations on input tensors across different nodes and then distributes the results to all nodes. For example, in model parallel training, the ReduceScatter operator is used to distribute the reduced gradients of partial model parameters to different computing nodes.

[0042] Specifically, the AllGather operator is an operator that collects data from all computing nodes and concatenates it into a complete result. For example, in data parallel training, the AllGather operator is used to collect local data on each GPU to form a complete global data view.

[0043] Specifically, the AlltoAll operator is an operator that performs full exchange of data between all computing nodes. For example, in hybrid parallel training, the AlltoAll operator is used to exchange intermediate calculation results between model slices of different dimensions.

[0044] In some embodiments, the memory access pattern in the collective communication operator can be identified by statically analyzing the compiler intermediate representation (IR) or analyzing the runtime program execution features.

[0045] Exemplarily, the memory access instruction sequence contained in the NVIDIA collective communications library (NCCL) source code or the compiled parallel thread execution (PTX) code can be identified.

[0046] Specifically, the memory access instructions in the kernel function can be extracted through the program analysis interface provided by the compiler, and the access pattern thereof can be analyzed.

[0047] Specifically, based on the low level virtual machine (LLVM) compilation framework, a custom analysis pass can be inserted in the compiler intermediate code optimization stage, which can identify the memory access pattern specific to the collective communication operator. By analyzing the operation code, operand type and address calculation pattern of the instruction, the occurrence position and access characteristics of the standard load instruction and the standard store instruction can be accurately identified.

[0048] Specifically, the specific memory access pattern can include continuous bulk memory access, data load and store pattern before and after the reduction operation, etc.

[0049] In this way, the present application accurately identifies the memory access instruction sequence in the collective communication operator, thereby providing an accurate optimization target for subsequent instruction level optimization, and ensuring the pertinence and effectiveness of the optimization operation.

[0050] S102, replace the standard load instruction with a batch load instruction.

[0051] In the embodiments of the present application, the batch load instruction is configured to batch load data from a continuous memory address space at the thread bundle level.

[0052] Optionally, the batch load instruction is an explicit batch operation instruction supported by hardware.

[0053] In some embodiments, the batch load instruction refers to an explicit batch memory read instruction supported at the GPU hardware level. For example, the ld.bulk instruction in the BR200 hardware architecture.

[0054] It should be noted that, unlike the traditional single-element load instruction, the batch load instruction can load multiple data elements in a continuous memory address space at a thread warp level at a time. For example, when a warp containing 32 threads executes the ld.bulk instruction, a 128-byte continuous data block can be loaded at a time, instead of 32 independent 4-byte load operations.

[0055] Exemplarily, in the BR200 GPU architecture, the ld.bulk instruction is implemented through a special hardware unit, which accepts a base address and a data size as parameters and can complete the loading operation of the continuous data required by the entire warp in a single instruction cycle.

[0056] It should be noted that, compared with the traditional ld.global instruction, the ld.bulk instruction reduces the number of instruction transmissions and improves the instruction-level parallelism.

[0057] In a specific implementation, this hardware function can be called through the built-in function __builtin_biren_ldbulk_##data_ptx_ty of the compiler. Wherein ##data_ptx_ty represents a data type placeholder defined in the PTX instruction set.

[0058] In this way, the application achieves the merged access to the continuous memory address space by replacing the standard load instruction with the batch load instruction at the thread warp level, significantly reduces the number of memory request transactions, and improves the utilization of memory bandwidth.

[0059] S103, replace the standard storage instruction with a batch storage instruction.

[0060] In the embodiments of the application, the batch storage instruction is configured to batch store data to a continuous memory address space at a thread warp level.

[0061] Optionally, the batch storage instruction is an explicit batch operation instruction supported by hardware.

[0062] In some embodiments, the batch storage instruction refers to an explicit batch memory write instruction supported at the GPU hardware level, such as the st.bulk instruction in the BR200 hardware architecture.

[0063] It should be noted that the batch storage instruction can write multiple data elements to a continuous memory address space at a warp level at a time. For example, when performing a gradient update operation, a warp can use a single st.bulk instruction to continuously write the updated 32 parameter values to the global memory, instead of transmitting 32 independent st.global instructions.

[0064] Exemplarily, in the BR200 GPU architecture, the st.bulk instruction realizes the batch write function through a special memory controller. The instruction supports configuring different storage properties, such as cache strategies and consistency protocols.

[0065] The L1 cache control attribute (such as NON_CACHEABLE) and the L2 cache control attribute (such as EVICT_NORMAL) can be set to optimize the storage performance in different use scenarios.

[0066] The L1 cache is usually shared by each streaming multiprocessor (SM), and the L2 cache is shared by all SMs.

[0067] In this way, the application replaces the standard storage instruction with the batch storage instruction at the thread bundle level, reduces the number of storage instruction launches, avoids storage conflicts that may occur in traditional storage operations, and improves the efficiency of storage operations.

[0068] S104, execute the batch load instruction and the batch storage instruction to reduce the number of memory access instruction launches and the scheduling burden of the thread bundle scheduler.

[0069] In some embodiments, when the batch load instruction is executed, the memory access address can be calculated by a scalar arithmetic logic unit.

[0070] Exemplarily, when batch loading is performed using the ld.bulk instruction, address calculation is completed by a scalar arithmetic logic unit (Scalar ALU) instead of using a traditional vector ALU. This design reduces the occupation of general integer pipeline (Integer Pipeline) resources.

[0071] Specifically, in the loop unrolling optimization scenario, the base address calculated by the scalar can be backed up in the thread bundle state register (warp state register, WSR) to avoid repeated address calculation in each loop iteration.

[0072] For example, when implementing gradient summation of the AllReduce operator, the base address of the gradient buffer can be calculated in advance and saved in the WSR, and the subsequent batch load instruction directly uses the address, reducing the number of address calculation instructions.

[0073] Therefore, the application effectively reduces the occupation of integer operation resources and improves the calculation efficiency by using a scalar ALU to perform address calculation and combining a register backup mechanism.

[0074] In some embodiments, when storing the instructions in batches, the data can be stored to the target memory address without accessing the staging buffer during the storage process.

[0075] For example, a conventional st.global instruction may need to first transmit data and address information to a staging buffer during execution, and then perform an actual storage operation by a memory controller. The st.bulk instruction directly stores data to a target memory address without passing through an intermediate staging buffer through hardware optimization.

[0076] Specifically, in the BR200 hardware architecture, the st.bulk instruction directly transmits data from a register to a target memory address through a dedicated storage path, reduces the data transmission link, and reduces the storage delay.

[0077] For example, after completing gradient reduction calculation, the st.bulk instruction is used to directly write the result to the global memory, avoiding the access overhead of the staging buffer and improving the utilization rate of the storage bandwidth.

[0078] Therefore, the application reduces the data transmission link, reduces the storage operation delay, and improves the utilization rate of the storage bandwidth by eliminating the access requirement of the staging buffer.

[0079] In the memory access optimization method for a collective communication operator provided by the application, the memory access instruction sequence in the collective communication operator is identified to accurately locate the optimization target; the standard load storage instruction is replaced by a batch operation instruction at the thread bundle level to combine originally dispersed single-element access into continuous batch access, effectively reducing the number of memory request transactions; the single emission of the batch instruction replaces the emission of multiple independent instructions, significantly reducing the instruction emission frequency and reducing the scheduling burden of the thread bundle scheduler; these steps cooperate with each other to convert the traditional fine-grained dispersed access mode into an efficient continuous block access mode, forming an optimization scheme that can efficiently and in batches access continuous memory blocks, thereby improving the execution efficiency and bandwidth utilization rate of GPU distributed training.

[0080] The performance advantages of the batch instruction are described below by taking the BR200 GPU as an example.

[0081] As shown in FIG. 8, it is a performance comparison diagram of the batch memory operation instruction based on the BR200. Figure 2

[0082] In the embodiments of the application, Figure 2 ​The horizontal axis represents the size of the contiguous memory data block to be processed, covering eight common data sizes in GPU distributed training: 256KB (kilobytes), 1MB (megabytes), 4MB, 16MB, 64MB, 128MB, 256MB, and 512MB. These data sizes cover the typical sizes of various tensors such as gradient data, weight data, and intermediate calculation results in large model training. The vertical axis represents the memory access performance metric, specifically memory access bandwidth, measured in gigabytes per second (GB / s). This metric directly reflects the speed of memory access; a higher value indicates better performance and is a core indicator for measuring memory operation efficiency.

[0083] For example, the graph contains 5 performance curves, each corresponding to a different optimized configuration and benchmark:

[0084] The first curve is "unroll4 (ld.bulk / st.bulk)", which represents the basic optimization configuration that uses batch load / store instructions and performs 4 iterations of code loop unrolling (unroll4).

[0085] The second curve is "unroll4-2GCU(ld.bulk / st.bulk)", which represents an optimized configuration that uses batch commands, loops four times, and two graphics compute units (GCUs).

[0086] The third curve is "unroll4-3GCU(ld.bulk / st.bulk)", which represents an optimized configuration that uses batch commands, loops four times, and three GCUs.

[0087] The fourth curve is "unroll4-4GCU(ld.bulk / st.bulk)", which represents an optimized configuration that uses batch commands, loops four times, and four GCUs.

[0088] The fifth curve is "NCCL-sim", which represents the benchmark reference group, that is, the performance level of the NCCL simulation environment (such as NCCL-sim) based on NVIDIA NV18 hardware (24 SMs) and unidirectional 450GB / s bandwidth, corresponding to the performance of traditional standard load / store instructions.

[0089] Specifically, from Figure 2 The performance data clearly shows that, across all test data scales, the implementation using batch memory operation instructions achieves higher effective bandwidth than the standard implementation.

[0090] For example, in a typical configuration (such as 2 GCUs), combined with loop unrolling optimization (such as unroll4), the performance of the batch instruction can reach 87.7% of the performance of the simulation competitor (such as NCCL) under similar conditions, and the performance bandwidth is improved by about 20% compared with the unoptimized basic version.

[0091] In addition, the optimized performance of different GCU numbers shows a stable trend, and the performance of the 2 GCU configuration is optimal in general, and the performance of the 3 and 4 GCU configurations is slightly lower than that of the 2 GCU configuration, but is better than the benchmark group.

[0092] Therefore, the above Figure 2 The intuitive verification of the batch memory operation instruction on the BR200 hardware verifies the optimization effect, and the quantitative data proves that the optimization scheme can improve the memory access bandwidth and solve the performance bottleneck of the traditional instruction in the processing of the continuous large data block.

[0093] Optionally, the batch load instruction and the batch store instruction can be called through the built-in function of the compiler.

[0094] For example, when optimizing in the CUDA compiler of the BR200 platform, the calling of the batch instruction can be realized by defining specific template functions and built-in functions.

[0095] For example, the following compiler built-in function interface can be defined:

[0096] c++

[0097] template<>

[0098] __device__ __forceinline__ BytePack <bytes>ld_##space <bytes>(addr_cxx_ty addr) {

[0099] auto gimaddr = (data_cxx_ty __attribute__((address_space(1))) ) addr;

[0100] BytePack <bytes>ans;

[0101] ans.native = __builtin_biren_ldbulk_##data_ptx_ty(gimaddr,

[0102] (int)_11_control_prop::EVICT_FIRST,

[0103] (int)_12_control_prop::EVICT_FIRST,

[0104] (int)supa::_access_attribute::rw, 0);

[0105] return ans;

[0106] }

[0107] Specifically, in the above code, __builtin_biren_ldbulk_##data_ptx_ty is a built-in function provided by the compiler for bulk loading, which accepts a target address parameter and multiple cache control parameters, and returns a byte package containing bulk data.

[0108] Similarly, for bulk storage operations, the __builtin_biren_stbulk_##data_ptx_ty built-in function can be used, which accepts a target address, data to be stored, and storage attribute parameters to complete the bulk storage operation.

[0109] In this way, the present application provides a convenient bulk instruction calling interface through the compiler built-in function, so that programmers can fully utilize the bulk operation capability of the hardware without understanding the underlying hardware details, improving the portability and development efficiency of the code.

[0110] In some embodiments, when calling the built-in function, parameters related to cache strategies can be configured.

[0111] In embodiments of the present application, the parameters include first-level cache control attributes and second-level cache control attributes.

[0112] Among them, the first-level cache control attribute refers to the parameter used to configure the L1 cache (level 1 cache) behavior, including non-cacheable mode (NON_CACHEABLE), priority replacement mode (EVICT_FIRST), etc.; the second-level cache control attribute refers to the parameter used to configure the L2 cache (level 2 cache) behavior, including priority replacement mode (EVICT_FIRST), normal replacement mode (EVICT_NORMAL), etc.

[0113] Exemplarily, when calling the __builtin_biren_ldbulk_##data_ptx_ty function, the L1 cache behavior can be set through the _11_control_prop parameter, and the L2 cache behavior can be set through the _12_control_prop parameter.

[0114] For example, for frequently accessed gradient data, the L1 cache can be set to EVICT_FIRST, and the L2 cache can be set to EVICT_NORMAL, to balance cache utilization and data consistency.

[0115] Specifically, different cache policy combinations can be adopted in different stages of distributed training. In the gradient aggregation stage, since the gradient data is usually used only once, the L1 cache can be set to NON_CACHEABLE to avoid polluting the cache; in the parameter update stage, since the parameters will be read multiple times, the L1 cache can be set to EVICT_FIRST, and the L2 cache can be set to EVICT_NORMAL to improve cache hit rate.

[0116] In this way, the application can optimize cache behavior for different computing scenarios through flexible cache policy configuration, and improve cache utilization and overall computing performance.

[0117] Optionally, the batch storage instruction supports burst transmission mode.

[0118] In the embodiment of the application, the batch storage instruction is used in the following scenarios by supporting burst transmission mode. For example: point-to-point data transmission scenario, multicast transmission scenario, atomic operation scenario.

[0119] Exemplarily, the st.bulk instruction supports st.bulk.lrbp.e4 variants, which support larger burst transmission length and can transmit more data in one transaction.

[0120] For example, in the peer-to-peer direct memory access (Direct P2P) scenario between GPUs, the batch storage instruction supporting burst transmission can be used to realize efficient inter-card data direct transmission.

[0121] Specifically, in the implementation of the AllReduce operator, when intermediate calculation results need to be exchanged between multiple GPUs, using the batch storage instruction supporting burst transmission can reduce the number of transactions and improve transmission efficiency.

[0122] For example, in the reduction-bulk mode, after each GPU completes the local reduction calculation, it uses the st.bulk.lrbp.e4 instruction to batch transfer the results to other GPUs. Compared with the traditional multiple store instructions, this reduces the number of bus transactions and improves the effective bandwidth.

[0123] Thus, by supporting batch storage instructions in burst transmission mode, this invention optimizes data transmission efficiency in point-to-point and multicast transmission scenarios, reduces bus transaction overhead, and improves overall system bandwidth.

[0124] Optionally, after obtaining the batch load instruction and the batch store instruction, the code containing the batch load instruction and the batch store instruction can be looped and expanded.

[0125] In this embodiment of the invention, loop unrolling refers to a compiler optimization technique that reduces the overhead of loop control instructions by copying the loop body multiple times.

[0126] In some embodiments, the thread bundle status register can be used to back up the scalar address in the loop control in order to complete the loop unrolling process.

[0127] For example, when implementing the gradient summation loop of the AllReduce operator, the loop can be unrolled four times (unroll4), and the base address pointer can be backed up using WSR in the unrolled code. This allows four consecutive blocks of data to be processed in each iteration, reducing the proportion of loop control instructions.

[0128] Specifically, in the CUDA code of the BR200 platform, optimization can be achieved by combining the compiler directive #pragma unroll 4 with manual address backup.

[0129] For example:

[0130] C++

[0131] #pragma unroll 4

[0132] for (int i = 0; i <element_count; i++) {

[0133] / / Use WSR to back up the base address

[0134] base_addr = __wsr_reg(base_ptr + i warp_size);

[0135] / / Use batch load command to read data

[0136] data = __builtin_biren_ldbulk_f32(base_addr,...);

[0137] / / Perform reduction computation

[0138] / / Write result using bulk store instruction

[0139] __builtin_biren_stbulk_f32(output_addr, result,...);

[0140] }

[0141] Thus, the application further reduces the loop control overhead and the number of address calculation instructions by combining the loop unrolling optimization with the register backup mechanism, improves the instruction level parallelism and the calculation efficiency, and makes the performance advantage of the bulk memory operation instruction fully play.

[0142] The memory access optimization device for the collective communication operator provided by the application is described below, and the memory access optimization device for the collective communication operator described below can be correspondingly referred to the memory access optimization method for the collective communication operator described above.

[0143] Figure 3 A structural schematic diagram of a memory access optimization device for a collective communication operator provided by an embodiment of the application. The memory access optimization device for the collective communication operator comprises an identification module 301, a load instruction replacement module 302, a storage instruction replacement module 303, and an execution module 304.

[0144] The identification module 301 is configured to identify a memory access instruction sequence in the collective communication operator to be optimized, and the instruction sequence comprises standard load instructions and standard storage instructions. The load instruction replacement module 302 is configured to replace the standard load instructions with bulk load instructions, and the bulk load instructions are configured to bulk load data from a continuous memory address space at a thread bundle level. The storage instruction replacement module 303 is configured to replace the standard storage instructions with bulk storage instructions, and the bulk storage instructions are configured to bulk store data to a continuous memory address space at a thread bundle level. The execution module 304 is configured to execute the bulk load instructions and the bulk storage instructions to reduce the number of memory access instructions and the scheduling burden of a thread bundle scheduler.

[0145] In some embodiments, the execution module 304 is specifically configured to calculate the memory access address by using a scalar arithmetic logic unit when executing the bulk load instructions.

[0146] In some embodiments, the execution module 304 is further configured to perform loop unrolling processing on a code loop containing the batch load instruction and the batch store instruction.

[0147] In some embodiments, the execution module 304 is specifically configured to back up the scalar address in the loop control using the thread bundle state register.

[0148] In the memory access optimization device for the collective communication operator provided by the application, the optimization target is accurately located by identifying the memory access instruction sequence in the collective communication operator; the originally dispersed single-element access is combined into continuous batch access by replacing the standard load store instruction with the batch operation instruction at the thread bundle level, thereby effectively reducing the number of memory request transactions; the instruction emission frequency is significantly reduced and the scheduling burden of the thread bundle scheduler is reduced by replacing multiple independent instruction emissions with a single batch instruction emission; these steps cooperate with each other to transform the traditional fine-grained dispersed access mode into an efficient continuous block access mode, forming an optimization scheme that can perform efficient batch access on continuous memory blocks, thereby improving the execution efficiency and bandwidth utilization of GPU distributed training.

[0149] Figure 4 An example of a physical structure diagram of a computing device is shown in FIG. 1. Figure 4 As shown in FIG. 1, the electronic device can include one or more graphic processors 410, a communications interface 420, a memory 430, and a communications bus 440, wherein the graphic processors 410, the communications interface 420, and the memory 430 communicate with each other through the communications bus 440. The graphic processors 410 can invoke the logical instructions in the memory 430 to execute the memory access optimization method for the collective communication operator, which includes: identifying the memory access instruction sequence in the collective communication operator to be optimized, the memory access instruction sequence including standard load instructions and standard store instructions; replacing the standard load instructions with batch load instructions and replacing the standard store instructions with batch store instructions; the batch load instructions and the batch store instructions are configured to perform batch operations on continuous memory address spaces at the thread bundle level; and finally executing the batch load instructions and the batch store instructions to reduce the number of memory access instruction emissions and reduce the scheduling burden of the thread bundle scheduler.

[0150] In addition, the logic instructions in the memory 430 described above can be implemented in the form of a software function unit and sold or used as an independent product, and can be stored in a computer readable storage medium. Based on such understanding, the technical solutions of the present application essentially or the parts that contribute to the prior art or 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 (read-only memory, ROM), a random access memory (random access memory, RAM), a magnetic disk or an optical disk, and various program code storage media.

[0151] 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 computer readable storage medium, and the computer program can be executed by a processor to enable a computer to execute the memory access optimization method for a collective communication operator provided in the above-mentioned embodiments. The method comprises: identifying a memory access instruction sequence in a collective communication operator to be optimized, the memory access instruction sequence comprising a standard load instruction and a standard store instruction; replacing the standard load instruction with a batch load instruction and replacing the standard store instruction with a batch store instruction; the batch load instruction and the batch store instruction are configured to perform batch operations on continuous memory address spaces at a thread bundle level; and finally executing the batch load instruction and the batch store instruction to reduce the number of memory access instruction transmissions and reduce the scheduling burden of a thread bundle scheduler.

[0152] In another aspect, the present application also provides a computer readable storage medium, which stores a computer program, the computer program is executed by a processor to implement the memory access optimization method for a collective communication operator provided in the above-mentioned embodiments. The method comprises: identifying a memory access instruction sequence in a collective communication operator to be optimized, the memory access instruction sequence comprising a standard load instruction and a standard store instruction; replacing the standard load instruction with a batch load instruction and replacing the standard store instruction with a batch store instruction; the batch load instruction and the batch store instruction are configured to perform batch operations on continuous memory address spaces at a thread bundle level; and finally executing the batch load instruction and the batch store instruction to reduce the number of memory access instruction transmissions and reduce the scheduling burden of a thread bundle scheduler.

[0153] It should be noted that the computer readable storage medium can be a non-transitory computer readable storage medium or a transitory computer readable storage medium, which is not limited here.

[0154] The device embodiments described above are merely illustrative, wherein the units described as separate components can or can not be physically separated, 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 purpose of the embodiment scheme according to actual needs. Those skilled in the art can understand and implement it without creative labor.

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

[0156] 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.< / bytes> < / bytes> < / bytes>

Claims

1. A memory access optimization method for collective communication operators, characterized in that, The method comprises: identifying a memory access instruction sequence in a collective communication operator to be optimized, the memory access instruction sequence comprising a standard load instruction and a standard store instruction, the standard load instruction being a memory read instruction executed independently by each thread, and the standard store instruction being a memory write instruction executed independently by each thread; replacing the standard load instruction with a bulk load instruction configured to bulk load data from a continuous memory address space at a thread bundle level; replacing the standard store instruction with a bulk store instruction configured to bulk store data to a continuous memory address space at a thread bundle level; executing the bulk load instruction and the bulk store instruction to reduce the number of memory access instruction launches and the scheduling burden of a thread bundle scheduler.

2. The method of claim 1, wherein, The collective communication operator comprises at least one of an all-reduce operator, a reduce-scatter operator, an all-gather operator, and an all-exchange operator.

3. The method of claim 1, wherein, The bulk load instruction and the bulk store instruction are hardware-supported explicit bulk operation instructions.

4. The method of claim 1, wherein, The bulk load instruction and the bulk store instruction are invoked by a built-in function of a compiler.

5. The method of claim 4, wherein, When the built-in function is invoked, a parameter related to a cache policy is configured, the parameter comprising a first-level cache control attribute and a second-level cache control attribute.

6. The method of claim 1, wherein, When the bulk load instruction is executed, a memory access address is calculated by a scalar arithmetic logic unit.

7. The method of claim 1, wherein, When the bulk store instruction is executed, data is stored to a target memory address without accessing a staging buffer during the storing process.

8. The method of claim 1, wherein, The bulk store instruction supports a burst transfer mode.

9. The method of claim 1, wherein, The method further comprises: performing loop unrolling processing on a code loop containing the bulk load instruction and the bulk store instruction.

10. The method of claim 9, wherein, The loop unrolling processing comprises: backing up a scalar address in a loop control by using a thread bundle state register.

11. A memory access optimization apparatus for collective communication operators, characterized in that, The apparatus comprises: an identifying module configured to identify a memory access instruction sequence in a collective communication operator to be optimized, the memory access instruction sequence comprising a standard load instruction and a standard store instruction, the standard load instruction being a memory read instruction executed independently by each thread, and the standard store instruction being a memory write instruction executed independently by each thread; a load instruction replacing module configured to replace the standard load instruction with a bulk load instruction configured to bulk load data from a continuous memory address space at a thread bundle level; a store instruction replacing module configured to replace the standard store instruction with a bulk store instruction configured to bulk store data to a continuous memory address space at a thread bundle level; an executing module configured to execute the bulk load instruction and the bulk store instruction to reduce the number of memory access instruction launches and the scheduling burden of a thread bundle scheduler.

12. A computing device, comprising: The computing device comprises: one or more graphics processors; a memory configured to store one or more programs, when the one or more programs are executed by the one or more processors, the one or more graphics processors are caused to implement the method according to any one of claims 1 to 10.

13. A computer readable storage medium having stored thereon a computer program, characterized in that, The computer program, when executed by a graphics processor, causes the graphics processor to perform the method of any one of claims 1 to 10.

Images