Gradient calculation method, device, storage medium and program product

Abstract

Embodiments of the present disclosure provide a gradient calculation method and device, a storage medium and a program product, and relate to the technical field of chips. The gradient calculation method comprises: determining a target gradient tensor that meets a preset gradient output condition from a candidate gradient tensor to be calculated; wherein the preset gradient output condition comprises that each output element of the target gradient tensor is generated by reduction operation based on a set reduction dimension, and different output elements are respectively written into different memory addresses of a memory of a target chip; and a plurality of target calculation tasks corresponding to a same output element in the target gradient tensor are allocated to different threads in a same thread bundle for execution, so that the thread bundle executes the plurality of target calculation tasks based on a register and a register interconnection network to obtain the output element after reduction operation of the set reduction dimension, and a specified thread in the thread bundle writes the output element into the memory of the target chip. In this way, the gradient calculation efficiency can be improved.

Description

Technical Field

[0001] This disclosure relates to the field of chip technology, and more specifically, to a gradient calculation method, apparatus, storage medium, and program product. Background Technology

[0002] In the training or inference of deep learning and other artificial intelligence (AI) models based on chips such as graphics processing units (GPUs), the computational efficiency of the backpropagation process has a significant impact on the training or inference efficiency of AI models. For example, in the backpropagation of deep learning operators, a large number of scalar gradients (such as attention weight gradients, sampling coordinate gradients, etc.) need to be reduced to specific dimensions (such as channel dimensions) before being written out.

[0003] In related technologies, the channel partial sums of each thread can be computed in parallel using multi-threaded GPUs, and then reduced in global memory using atomic operations such as atomic addition. This results in multiple atomic operations being triggered for each gradient calculation. Since atomic operations must pass serially through the arbitration queue of the GPU cache controller, access latency and memory bandwidth contention occur, affecting gradient computation efficiency. Summary of the Invention

[0004] One objective of this disclosure is to provide a new technical solution for gradient calculation.

[0005] According to a first aspect of the present disclosure, a gradient calculation method is provided, applied to a target chip, the target chip including multiple thread bundles, each thread bundle including multiple threads, a register corresponding to each thread, and a register interconnection network for interconnecting access of multiple registers; the method includes: From the candidate gradient tensors to be calculated, a target gradient tensor that satisfies the preset gradient output conditions is determined; wherein, the preset gradient output conditions include: each of the multiple output elements of the target gradient tensor is generated by a reduction operation based on a set reduction dimension, and different output elements are written to different storage addresses in the memory of the target chip. Multiple target computation tasks corresponding to the same output element in the target gradient tensor are assigned to different threads within the same thread bundle for execution. This allows the thread bundle to execute multiple target computation tasks based on the register and the register interconnection network to perform a reduction operation on the set reduction dimension and obtain the output element. The designated thread within the thread bundle then writes the output element into the memory of the target chip.

[0006] Optionally, the element representation corresponding to the output element has multiple dimensions; the step of allocating multiple target computation tasks corresponding to the same output element in the target gradient tensor to different threads within the same thread bundle for execution includes: Convert the multiple element representation dimensions of the output element into a one-dimensional element index corresponding to the output element; The thread bundle identifier corresponding to the output element is determined based on the one-dimensional element index; For any output element, the multiple target computation tasks corresponding to the output element are assigned to different threads within the thread bundle corresponding to the thread bundle identifier for execution.

[0007] Optionally, converting the multiple element representation dimensions of the output element into a one-dimensional element index corresponding to the output element includes: Determine the dimension index value of the output element in each of the dimensions represented by the element; Obtain the number of dimension indices of the multiple output elements of the target gradient tensor in each of the dimension representations of the elements; Based on the dimension index value and the number of dimension indices, the one-dimensional element index corresponding to the output element is calculated.

[0008] Optionally, in the multiple target computation tasks corresponding to the same output element, each target computation task corresponds one-to-one with the reduction dimension index value corresponding to the set reduction dimension of the output element; the step of allocating the multiple target computation tasks corresponding to the same output element in the target gradient tensor to different threads within the same thread bundle for execution includes: Determine the number of threads within the thread bundle and the number of reduction dimension indices for the set reduction dimension; Based on the number of threads and the number of reduction dimension indices, determine the correspondence between the thread index value of each thread within the thread bundle and the thread allocation index value of the reduction dimension index; Based on the thread allocation index correspondence, the target computation tasks corresponding to different reduction dimension index values ​​are allocated to the target threads corresponding to the reduction dimension index values.

[0009] Optionally, the candidate gradient tensor is the gradient tensor within the target operator, and determining the target gradient tensor that satisfies the preset gradient output condition from the candidate gradient tensors to be calculated includes: During the operator development phase, candidate gradient tensors to be computed in the target operator are determined; From the candidate gradient tensors to be calculated, determine the candidate gradient tensors that satisfy the preset gradient output conditions as the target gradient tensor; The method further includes: taking candidate gradient tensors that do not meet the preset gradient output conditions as other gradient tensors; wherein the gradient calculation of the other gradient tensors adopts a different task allocation method than that of the target gradient tensor.

[0010] Optionally, the preset gradient output condition further includes: the output dimension of the output element does not include the set reduction dimension.

[0011] Optionally, the thread bundle can implement the reduction operation for multiple channels in the following ways: Each thread in the thread bundle executes the assigned target computation task based on the register corresponding to the thread to calculate the thread-level reduction result of at least one channel, and outputs the thread-level reduction result to the register; The reduction operation is performed on multiple thread-level reduction results within the thread bundle using a register interconnect network to obtain the element-wise reduction result of the output element.

[0012] Optionally, the step of performing reduction operations on multiple thread-level reduction results within the thread bundle via a register interconnect network to obtain the output elements of the target gradient tensor includes: The thread bundle shuffle instruction is executed within the thread bundle through the register interconnection network to perform tree reduction on multiple thread-level reduction results, thereby obtaining the element reduction result of the output element.

[0013] Optionally, the target gradient tensor includes weight gradients and / or coordinate gradients.

[0014] Optionally, the reduction dimension setting includes a channel dimension, and the number of reduction dimension indices of the channel dimension is an integer multiple of the number of threads within the thread bundle.

[0015] According to a second aspect of the present disclosure, an electronic device is provided, including a memory and a processor, the memory being configured to store computer instructions, and the processor being configured to invoke the computer instructions from the memory to perform the method described in the first aspect.

[0016] According to a third aspect of the present disclosure, a computer-readable storage medium is provided having a computer program stored thereon, the computer program implementing the method described in the first aspect when executed by a processor.

[0017] According to a fourth aspect of the present disclosure, a computer program product is provided, the computer program product including a computer program that, when executed by a processor, implements the method described in the first aspect.

[0018] Based on the gradient calculation method provided in this disclosure, the reduction operation of the same output element is assigned to a thread bundle, thereby transferring the reduction operation of the target gradient tensor from the traditional atomic accumulation path based on the global storage hierarchy to the register interconnection network within the thread bundle. This can reduce or eliminate the occupation of the cache controller CAS queue by related atomic operations, reduce storage access latency, alleviate storage bandwidth contention, and realize hardware acceleration of the reduction operation by utilizing extremely low latency registers and register interconnection networks, thereby improving the gradient calculation efficiency.

[0019] Other features and advantages of this disclosure will become clear from the following detailed description of exemplary embodiments with reference to the accompanying drawings. Attached Figure Description

[0020] The accompanying drawings, which are incorporated in and form part of this specification, illustrate embodiments of the present disclosure and, together with the description, serve to explain the principles of the disclosure.

[0021] Figure 1 This is a schematic diagram of an application scenario in which the technical solutions provided in the embodiments of this disclosure can be applied.

[0022] Figure 2 This is a flowchart illustrating a gradient calculation method provided in an embodiment of this disclosure.

[0023] Figure 3 This is a schematic diagram of the architecture of a gradient calculation system provided in an embodiment of this disclosure.

[0024] Figure 4 This is a schematic diagram of the structure of an electronic device provided in an embodiment of this disclosure. Detailed Implementation

[0025] Various exemplary embodiments of this disclosure will now be described in detail with reference to the accompanying drawings.

[0026] The following description of at least one exemplary embodiment is merely illustrative and is not intended to limit this disclosure or its application or use. Techniques, methods, and apparatus known to those skilled in the art may not be discussed in detail, but where appropriate, such techniques, methods, and apparatus may be considered part of the specification.

[0027] It should be noted that similar labels and letters in the following figures indicate similar items, so once an item is defined in one figure, it need not be discussed further in subsequent figures.

[0028] This disclosure relates to a technical solution for gradient calculation. Figure 1 This is a schematic diagram illustrating an application scenario where the technical solutions provided in the embodiments of this disclosure can be applied. For example... Figure 1 As shown, this application scenario may include a gradient calculation platform 100 and an operation terminal 200.

[0029] The gradient calculation platform 100 can be used to perform gradient calculations, such as executing the gradient calculation method provided in this disclosure embodiment. This gradient calculation method can be triggered by an AI agent and / or human engineer deployed on the gradient calculation platform. The gradient calculation platform may include a target chip 101, on which the gradient calculation method is executed. The target chip can be one or more chips of any type, such as a graphics processing unit (GPU), a general-purpose graphics processing unit (GPGPU), a neural network processing unit (NPU), a tensor processing unit (TPU), or a central processing unit (CPU). For example, the target chip 101 can be a chip with a Single Instruction Multiple Threads (SIMT) architecture. In some examples, the target chip may include multiple warp threads, each warp thread including multiple lines, a register corresponding to each line, and a register interconnection network for interconnecting access of multiple registers. The target chip may execute target computation tasks based on multiple lines in the warp thread to implement some or all of the steps of the gradient computation method provided in the embodiments of this disclosure.

[0030] In some examples, the gradient calculation platform may further include a memory, which may include at least one of static random access memory (SRAM), read-only memory (ROM), or erasable programmable read-only memory (EPROM). The memory can be used to store a computer program that, when executed by a processor, controls the processor to perform some or all of the steps of the gradient calculation method of the embodiments of this disclosure.

[0031] The operating terminal 200 and the gradient calculation platform 100 can establish a connection via an interface device or a communication device. Users can transmit gradient calculation-related data, files, or other information (such as operation instructions) to the gradient calculation platform via the operating terminal. The gradient calculation platform can read the user-input data, files, or other information and execute the gradient calculation method of this embodiment. Users can also manage the execution process of the gradient calculation method on the gradient calculation platform via the operating terminal; for example, users can start, pause, stop, or exit the execution process of the gradient calculation method.

[0032] exist Figure 1 In the application scenario shown, the gradient calculation platform 100 and the operation terminal 200 are illustrated as different physical devices, and the target chip 101 is illustrated as part of the gradient calculation platform 100. It should be noted that the gradient calculation platform 100 and the operation terminal 200 can be integrated on the same physical device, and the target chip 101 and the gradient calculation platform 100 can also be two different physical devices. This disclosure does not limit this aspect.

[0033] by Figure 1 Taking the illustrated application scenario as an example, in the backpropagation of deep learning operators based on the target chip, a large number of scalar gradients need to be reduced to specific dimensions (e.g., channel dimensions) before being written out. This reduction operation can include reduction summation, reduction to maximum value, reduction to minimum value, or reduction to average value. For example, attention weight gradients and sampling coordinate gradients need to be reduced and summed to the channel dimension before being written out. In related technologies, each thread can process data for one channel dimension, calculate the partial derivative contribution of that channel, and then accumulate the result to the target address in global memory using atomic operations such as `atomicAdd`. For a typical configuration of 32 channel dimensions, each output element can generate 32 atomic operations such as `atomicAdd`. Since atomic operations must pass serially through the arbitration queue of the GPU cache controller, access latency and memory bandwidth contention will occur, affecting gradient calculation efficiency.

[0034] Taking operators involving spatial sampling and weighted aggregation patterns (such as deformable convolution, deformable attention, and region of interest alignment) as examples, the backpropagation of these operators requires the computation of multiple sets of gradients. Each output element of the weight gradient *w* and the coordinate gradient *loc* corresponds to a unique storage address, and the write addresses between different output elements are completely different, eliminating write conflicts. However, in related technologies, these conflict-free gradients and other gradients with write conflicts (such as eigenvalue gradients) are all processed using atomic operations. Each conflict-free gradient output consumes multiple atomic operations, and these atomic operations are processed through a Compare-And-Swap (CAS) queue path in the cache controller. Even if there is no actual contention, each CAS still requires additional comparison and swap cycles and queue management overhead, making it much slower than ordinary storage instructions.

[0035] In related technologies, all gradients are coupled into a single computational kernel function, and atomic operations are used uniformly for all gradients. In reality, the write conflict patterns of different gradient tensors are quite different: weight gradients and coordinate gradients have no write conflict, while eigenvalue gradients have a high write conflict (high probability of write conflict), making it impossible to match the optimal hardware primitive for each set of gradients, resulting in a waste of hardware resources.

[0036] Furthermore, GPU hardware provides an intra-thread register interconnect network, accessible via shuffle instructions. This network has a latency of only 1 to 2 clock cycles, bandwidth equivalent to register reads, and does not occupy load or storage units. However, related technologies completely fail to utilize this hardware resource for channel-level reduction operations. Instead, they accumulate gradients atomically at a global storage level with latency two orders of magnitude higher, resulting in low gradient computation efficiency.

[0037] To address the problems in related technologies, this disclosure provides a gradient calculation method that transfers the reduction operation of the target gradient tensor that meets the preset gradient output conditions from the global storage level to the register interconnection network within the thread bundle. The reduction result (i.e., the output element obtained by reduction) is written out by a designated thread within the thread bundle. For example, the reduction result can be written out atomically, thereby reducing or eliminating related atomic operations, reducing the processing pressure of the cache controller's CAS queue, and helping to improve the efficiency of gradient calculation.

[0038] Figure 2 This is a flowchart illustrating a gradient calculation method provided in an embodiment of this disclosure. This gradient calculation method can be applied to... Figure 1 The gradient computation platform shown can be applied, for example, to a target chip within a gradient computation platform. This target chip may include multiple thread bundles, each thread bundle comprising multiple threads, a register corresponding to each thread, and a register interconnect network for interconnecting and accessing multiple registers. For example... Figure 2As shown, the gradient calculation method in this embodiment may include the following steps S210 to S220.

[0039] Step S210: Determine the target gradient tensor that satisfies the preset gradient output conditions from the candidate gradient tensors to be calculated.

[0040] In some examples, the candidate gradient tensors to be computed can be all gradient tensors within the target operator, which can be any operator used in the artificial intelligence model. For example, for operators that include spatial sampling and weighted aggregation patterns, such as deformable convolution, deformable attention, region of interest (RoI) alignment, and grid sampling, backpropagation requires computing multiple sets of gradient tensors, including weight gradients, coordinate gradients, and eigenvalue gradients, all of which can serve as candidate gradient tensors.

[0041] In some examples, the preset gradient output condition may include: each of the multiple output elements of the target gradient tensor is generated by a reduction operation based on a set reduction dimension, and different output elements are written to different memory addresses in the target chip's memory. The set reduction dimension can be the dimension on which the reduction operation is performed on the output elements, such as the channel dimension or other reduction dimensions.

[0042] For example, the preset gradient output condition may further include: the output dimension of the output element does not include the set reduction dimension. For example, taking the weight gradient as an example, the set reduction dimension of the weight gradient can be the channel dimension, and the output dimension of the weight gradient may not include the channel dimension. Each output element can be a scalar result obtained after complete reduction of the channel dimension, and different output elements can be written to different storage addresses without write conflicts. This can exclude gradient tensors that still retain the set reduction dimension in the output dimension, ensuring that each output element of the target gradient tensor is a scalar result obtained after complete reduction of the set reduction dimension, avoiding overlapping write addresses between different output elements, and improving the reliability of gradient calculation.

[0043] For example, the preset gradient output condition may also include: different threads within the same thread bundle are responsible for calculating the components of the same output element at different reduction dimension indices. For example, the reduction dimension is set as the channel dimension, and the components at the reduction dimension index can be channel components.

[0044] In some examples, the preset gradient output conditions may include one or more of the following: Each of the multiple output elements of the target gradient tensor is generated by a reduction operation based on a set reduction dimension. Different output elements are written to different memory addresses in the target chip's memory; The output dimension of the output element does not include the reduced dimension specified therein; Within the same thread bundle, different threads are responsible for calculating the components of the same output element on different reduction dimension indices.

[0045] In this way, based on one or more of the above-mentioned preset gradient output conditions, gradient tensors that are not suitable for reduction within a thread bundle can be filtered out in advance during the task allocation stage, which helps to reduce the waste of computing resources caused by misjudgment and further improve the accuracy and efficiency of gradient calculation.

[0046] In some examples, the target gradient tensor among the candidate gradient tensors that satisfies the preset gradient output conditions can be called a first-type gradient tensor or a type A gradient tensor. This first-type gradient tensor can be a channel-reduced, write-conflict-free gradient tensor, for example, where each output element is the sum of the channel dimensions, and different output elements are written to different addresses. In some examples, the target gradient tensor may include weight gradients (w), coordinate gradients (loc), or other gradient tensors that satisfy the preset gradient output conditions.

[0047] In other examples, other gradient tensors among the candidate gradient tensors that do not satisfy the preset gradient output conditions can be called second-type gradient tensors or type B gradient tensors. These second-type gradient tensors can be spatially scattering gradient tensors with write conflicts, such as multiple computation sources writing to the same target address. These second-type gradient tensors can include eigenvalue gradients (V) or other gradient tensors that do not satisfy the aforementioned preset gradient output conditions.

[0048] The gradient computation of other gradient tensors can be handled using a different task allocation method than that of the target gradient tensor. For example, the target gradient tensor and other gradient tensors can be assigned to different computation kernels for execution, thus achieving kernel decomposition. In this way, optimal hardware primitives can be matched to gradient tensors with different conflict modes, helping to avoid wasting atomic operation resources in write-conflict scenarios and improving gradient computation efficiency.

[0049] In some examples, during the operator development phase, candidate gradient tensors to be computed in the target operator can be identified. From these candidate gradient tensors, those satisfying preset gradient output conditions can be selected as the target gradient tensor. Alternatively, candidate gradient tensors that do not satisfy the preset gradient output conditions can be used as other gradient tensors, whose gradient computation can employ different task allocation methods than the target gradient tensor. For instance, static analysis of candidate gradient tensors can be performed during the operator development phase to determine the target gradient tensor satisfying the preset gradient output conditions. Exemplarily, the conflict category of each gradient tensor can be determined by analyzing the write address mapping relationship of each candidate gradient tensor. This write address mapping relationship, also known as the write conflict topology, can be used to describe the conflict relationship between the write target addresses of multiple threads in parallel computing. For example, if two threads write to the same address, there is a conflict, and the preset gradient output condition is not met; if the write addresses are different, there is no conflict, and the preset gradient output condition is met. By analyzing the write conflict topology, gradient tensors can be divided into two categories: those without write conflicts and those with write conflicts. The target gradient tensor satisfying the preset gradient output condition can be a candidate gradient tensor without write conflicts.

[0050] For example, taking deformable convolution, deformable attention, and region-of-interest alignment as examples, the backpropagation of these operators requires calculating multiple sets of gradients. Among them, the weight gradient (also known as the attention weight gradient) w) and coordinate gradient (also known as sampled coordinate gradient) `loc` satisfies the preset gradient output conditions and belongs to the target gradient tensor. Taking the weighted gradient as an example, each output element is a scalar value obtained after reducing the channel dimension. Different output elements correspond to different storage addresses, and there is no write conflict. Eigenvalue gradient ( The output dimension of V includes the channel dimension, and multiple computation sources may write to the same target address, resulting in write conflicts. It belongs to other gradient tensors and can be handled using different task allocation methods.

[0051] In this way, static analysis during the operator development phase can pre-define gradient classification results, eliminating the need for repeated judgments at runtime and helping to reduce runtime overhead. Assigning gradients with different conflict modes to different kernel functions allows for matching the optimal hardware primitives for each set of gradients, helping to avoid wasting atomic operation resources in write-free conflict scenarios and improving gradient computation efficiency.

[0052] In other examples, the target gradient tensor can be dynamically determined at runtime based on a judgment logic that predetermines the gradient output conditions. For instance, before operator execution, it can be checked whether the output dimension of the gradient tensor includes the set reduction dimension and whether different output elements are written to different storage addresses; if the conditions are met, it is determined as the target gradient tensor. Dynamic determination at runtime can adapt to different operator configurations, helping to improve the flexibility of the method.

[0053] In some examples, the number of reduced dimension indices can be greater than or equal to the number of threads within the thread bundle. For instance, the reduced dimension can be a channel dimension, and the number of reduced dimension indices (also called the number of channels C) can be greater than or equal to the number of threads within the thread bundle. For example, the number of reduced dimension indices can be an integer multiple of the number of threads within the thread bundle. For example, the number of threads within the thread bundle can be 32, and the number of channels C can be 32, 64, 128, etc., meaning the number of channels C can be greater than or equal to the number of threads within the thread bundle. This fully utilizes the parallel computing power within the thread bundle, further improving gradient calculation efficiency.

[0054] Step S220: Multiple target computation tasks corresponding to the same output element in the target gradient tensor are assigned to different threads within the same thread bundle for execution, so that the thread bundle executes multiple target computation tasks based on registers and register interconnection networks to achieve the reduction operation of the set reduction dimension and obtain the output element. The specified thread within the thread bundle writes the output element into the memory of the target chip.

[0055] For example, a target gradient tensor may include one or more output elements, each output element may correspond to multiple target computation tasks, and each target computation task may correspond to a reduction dimension index value that sets the reduction dimension. For example, if the reduction dimension is set to a channel dimension, then each target computation task may correspond to a channel index value.

[0056] In some examples, different threads within a thread bundle can execute target computation tasks corresponding to different reduction dimension index values. Each thread can execute its assigned target computation task based on its corresponding register (i.e., the thread-in-thread register) to calculate the thread-level reduction result. It should be noted that a thread can be assigned to execute one target computation task (i.e., corresponding to one reduction dimension index value) or multiple target computation tasks (i.e., corresponding to multiple reduction dimension index values). This embodiment does not limit this.

[0057] In some examples, a thread bundle can perform reduction operations on the reduction results of multiple threads within the bundle through a register interconnect network to obtain output elements (e.g., element-wise reduction results of output elements). The register interconnect network can be an on-chip direct interconnect path between the register files of multiple threads within a thread bundle on the target chip. Through the register interconnect network, different threads within a thread bundle can directly access data in the registers of other threads without going through global or shared memory. The register interconnect network can be accessed via the thread bundle shuffle instruction, with an access latency of 1 to 2 clock cycles, significantly lower than the access latency of global or shared memory.

[0058] A designated thread within a thread bundle can write the output elements to the target chip's memory. For example, this designated thread could be thread number 0 within the thread bundle, also known as lane 0. The operation of writing the element reduction result of the output elements to the target chip's memory by this designated thread can be either atomic or non-atomic. For instance, it can be written to global memory via a non-atomic direct write instruction. This non-atomic direct write instruction can be written via a cached write-back path, without comparison and operation queue swapping, thus improving write efficiency.

[0059] Furthermore, other threads within this thread bundle, besides the designated thread, can skip writing to the target chip's memory and instead run the target computation task through registers, writing the obtained thread-level reduction result into the registers, thereby improving efficiency.

[0060] Using the methods described in steps S210 to S220, target gradient tensors that meet preset gradient output conditions are identified from candidate gradient tensors. Multiple target computation tasks corresponding to the same output element are centrally allocated to different threads within the same thread bundle for execution. The reduction operation with the set reduction dimension is completed using registers and register interconnects within the thread bundle, and the output element is written to memory in a non-atomic manner by a designated thread. In this way, the reduction operation of the same output element is allocated to a thread bundle, thereby shifting the reduction operation of the target gradient tensor from the traditional atomic accumulation path based on the global storage hierarchy to the register interconnects within the thread bundle. This reduces or eliminates the occupation of the cache controller's CAS queue by related atomic operations, reduces storage access latency, alleviates storage bandwidth contention, and achieves hardware acceleration of the reduction operation using extremely low-latency registers and register interconnects, thus improving gradient computation efficiency.

[0061] In some embodiments of this disclosure, the element representing the above-mentioned output element has multiple dimensions, which can be the input dimension and / or the output dimension.

[0062] In some examples, the element representation dimension corresponding to the output element may not include the aforementioned reduced dimension, or the output dimension may not include the aforementioned reduced dimension, but the input dimension may include the reduced dimension. Taking the reduced dimension as a channel dimension as an example, the element representation dimension corresponding to the output element can be a non-channel dimension. For example, taking a deformable attention operator, the element representation dimension of the output element may include multiple element representation dimensions such as batch dimension, query position dimension, and attention head dimension.

[0063] For example, the method of assigning multiple target computation tasks corresponding to the same output element in the target gradient tensor to different threads within the same thread bundle in step S220 above may include: converting the multiple element representation dimensions of the output element into a one-dimensional element index corresponding to the output element, and determining the thread bundle identifier corresponding to the output element based on the one-dimensional element index; different output elements can be assigned to different thread bundles corresponding to different thread bundle identifiers. For example, for any output element, multiple target computation tasks corresponding to the output element can be assigned to different threads within the thread bundle corresponding to the thread bundle identifier.

[0064] For example, different output elements in the target gradient tensor can correspond to different one-dimensional element indices, and different one-dimensional element indices can also correspond to different output elements. In this way, an output element can be uniquely identified through a one-dimensional element index, and further, a thread bundle can be uniquely identified.

[0065] In this example, the dimension index value of the output element in each element representation dimension can be determined, and multiple dimension index values ​​can be converted into one-dimensional element indices. Specific conversion methods can include linear mapping conversion, hash function mapping conversion, or other mapping conversion methods that meet the uniqueness requirement.

[0066] In some examples, the linear mapping transformation can include: determining the dimensional index value of the output element in each element representation dimension; obtaining the number of dimensional indices of the multiple output elements of the target gradient tensor in each element representation dimension; and calculating the one-dimensional element index corresponding to the output element based on the dimensional index value and the number of dimensional indices. For example, the one-dimensional element index can be calculated through linear mapping. In this way, multi-dimensional indices are compressed and mapped to one-dimensional element indices through linear mapping, which is computationally simple and efficient and helps to reduce the computational overhead of index transformation.

[0067] Each output element of the target gradient tensor can have N element-wise representation dimensions, where N is an integer greater than 1. Taking an output element with three element-wise representation dimensions (i.e., N=3) as an example, each output element has a dimension index value in each element-wise representation dimension. For example, the dimension index value of an output element can be represented as (b, q, h), where b is the dimension index value of the batch dimension, q is the dimension index value of the query position dimension, and h is the dimension index value of the attention head dimension. Furthermore, multiple output elements of this target gradient tensor can have a number of dimension indices in each element-wise representation dimension. For example, the number of dimension indices for the batch dimension can be B, for the query position dimension can be Q, and for the attention head dimension can be H. Based on the dimension index value and the number of dimension indices, the one-dimensional element index corresponding to the output element can be calculated. This one-dimensional element index can correspond one-to-one with the warp ID, for example, the two can be equal.

[0068] In one implementation, a linear mapping can be performed in the order of batch dimension, query position dimension, and attention head dimension. The one-dimensional element index is calculated as (b×Q×H + q×H + h), that is, with the batch dimension as the highest bit and the attention head dimension as the lowest bit, it is expanded according to the product relationship of the dimensions to obtain a unique one-dimensional number. In another implementation, the order of the dimensions can be adjusted. For example, with the attention head dimension as the highest bit, the one-dimensional element index is calculated as h×B×Q + b×Q + q. That is, it is first expanded on the attention head dimension, and then expanded in sequence on the batch dimension and query position dimension, which also yields a unique one-dimensional number. Both methods can ensure that the one-dimensional element indices corresponding to different output elements are different, which is used to determine the thread bundle identifier later.

[0069] In other examples, the hash function mapping transformation may include: determining the dimension index value of the output element in each element representation dimension, inputting the dimension index value into a pre-designed hash function, and calculating the one-dimensional element index. For example, inputting the multi-dimensional index (b, q, h) into the hash function calculates the one-dimensional element index. The hash function can be designed to output different hash values ​​for any different multi-dimensional index input, thereby making the one-dimensional element index unique. Exemplarily, the hash function can be constructed based on a weighted XOR operation of the dimension index values ​​to achieve a unique mapping, or the hash function can be implemented in other ways in related technologies, which are not limited in this embodiment.

[0070] By converting multidimensional indices into one-dimensional thread bundle identifiers, channel computation tasks corresponding to the same multidimensional index can be concentrated in the same thread bundle, providing hardware conditions for subsequent reduction operations to be completed within the register interconnect network, thereby improving the efficiency of reduction operations.

[0071] In some examples, each target computation task corresponding to the same output element corresponds one-to-one with the reduction dimension index value of the output element. For example, if the reduction dimension is set to a channel dimension, then each target computation task can correspond to a channel index value c.

[0072] After determining the thread bundle corresponding to the output element, the method in step S220 above of assigning multiple target computation tasks corresponding to the same output element in the target gradient tensor to different threads within the same thread bundle for execution can further include: determining the number of threads within the thread bundle and setting the number of reduction dimension indices; determining the thread allocation index correspondence between the thread index value of each thread within the thread bundle and the reduction dimension index value based on the number of threads and the number of reduction dimension indices; and assigning the target computation tasks corresponding to different reduction dimension index values ​​to the target threads corresponding to the reduction dimension index values ​​based on the thread allocation index correspondence. This thread allocation index correspondence can include a correspondence between thread index values ​​and reduction dimension index values, and can be a one-to-many relationship.

[0073] For example, the number of threads W within a thread bundle and the number of reduction dimension indices C can be determined. For instance, the number of threads W can be 32, and the number of reduction dimension indices C can be 64. In an alternative implementation, the thread allocation index correspondence can adopt a staggered approach. For example, a thread with a thread index value of k can be assigned target computation tasks with reduction dimension indices of k, k+W, k+2W, etc. For example, if the number of threads W is 32 and the number of reduction dimension indices C is 64, then the thread numbered 0 can be responsible for target computation tasks with channel indices of 0 and 32, the thread numbered 1 can be responsible for target computation tasks with channel indices of 1 and 33, and so on.

[0074] It should be noted that other methods can also be used to determine the thread allocation index correspondence. For example, if the number of reduced dimension indexes C is M times the number of threads W, then the C reduced dimension indexes are divided into W parts, each containing M consecutive reduced dimension indexes, and the M consecutive reduced dimension indexes are allocated to one thread index.

[0075] In some embodiments of this disclosure, the method by which a thread bundle implements reduction operations for multiple channels may include: each thread of the thread bundle can execute an assigned target computation task based on the register corresponding to that thread to calculate a thread-level reduction result for at least one channel, and output the thread-level reduction result to a register; and perform reduction operations on multiple thread-level reduction results within the thread bundle through a register interconnection network to obtain an element-wise reduction result for the output elements. Exemplarily, this thread-level reduction result may also be referred to as a channel partial sum, and the element-wise reduction result may also be referred to as a full-thread bundle reduction result or a global sum.

[0076] For example, taking the weight gradient of the deformable attention operator as an example, the thread-level reduction result (channel partial sum) corresponding to thread k can be calculated as follows: reducing (e.g., summing) the channel components corresponding to all channel index values ​​under the responsibility of this thread. This thread can store the thread-level reduction result in a private register.

[0077] In this way, the thread-level reduction result is stored in a register. Register interconnect reduction does not generate memory access requests and does not occupy load or store functional units, which helps reduce memory access latency and improve reduction efficiency. It should be noted that the thread-level reduction result can also be stored in shared memory for reduction operations. However, the access latency of shared memory is higher than that of register interconnect.

[0078] In some examples, the method of obtaining the output elements of the target gradient tensor by performing reduction operations on multiple thread-level reduction results within a thread bundle through a register interconnect network can include: executing a thread bundle shuffling instruction within the thread bundle through the register interconnect network to perform tree-like reduction on multiple thread-level reduction results, thereby obtaining the element-wise reduction result of the output element. This thread bundle shuffling instruction can include an XOR shuffling synchronization instruction (e.g., the __shfl_xor_sync instruction).

[0079] For example, if a thread bundle includes 32 threads, five rounds of XOR operations can be performed. In the first round, the step size can be 16, and the thread numbered k can exchange register data with the thread numbered k+16 through the register interconnect network and accumulate the data, achieving data merging across half of the thread bundle. For example, the thread numbered 0 can exchange data with the thread numbered 16 and accumulate the data, the thread numbered 1 can exchange data with the thread numbered 17 and accumulate the data, and so on. In the second round, the step size can be reduced to 8, and the thread numbered k can exchange data with the thread numbered k+8 and accumulate the data. In the third round, the step size can be reduced to 4. In the fourth round, the step size can be reduced to 2. In the fifth round, the step size can be reduced to 1. After five rounds of operations, all threads within the thread bundle can hold the same complete reduction result.

[0080] For example, this tree reduction process can be represented as: S(r) = S(r-1) + shfl_xor(S(r-1), 2 （5 -r） ), where r can be the round number, r can be 1 to 5, shfl_xor is the XOR shuffle operation, S(r) represents the reduction result of the r-th round, and S(r-1) represents the reduction result of the (r-1)-th round. After the r-th round of reduction, every 2 r One consecutive thread can hold the same partial reduction result (e.g., partial sum or channel partial sum), and after 5 rounds, all 32 threads can hold the same complete reduction result (e.g., global sum).

[0081] Exemplarily, the reduction operation can be performed within the register interconnect network without generating any storage access requests or occupying load or store functional units (e.g., load / store units). Exemplarily, the latency of each round of XOR shuffling operation can be 1 to 2 clock cycles, and the total latency for 5 rounds can not exceed 10 clock cycles. Compared to the latency of atomic operations passing through the CAS queue in related technologies (exemplarily, each latency can be 50 to 200 clock cycles), the reduction method of this disclosure can significantly reduce latency and help improve gradient calculation efficiency. For a typical configuration with 32 channel indices, the calculation of each output element can generate 32 atomic operations. However, through the method of this disclosure, these 32 atomic operations can be reduced, eliminating the processing pressure of the L2 cache controller CAS queue on this part of the gradient and freeing up storage bandwidth. In addition, the thread bundle shuffling instruction can use an independent register interconnect functional unit, without occupying the issue slot of the load or store unit, which can avoid competition for issue ports with feature value readings and help improve instruction issue efficiency.

[0082] It should be noted that reduction operations within a thread bundle can also be performed using other reduction methods based on register interconnect networks, such as shift-shuffle reduction or other tree-structured reduction methods. Other reduction topologies, such as butterfly reduction or fan-shaped reduction, can also be employed.

[0083] In some examples, the aforementioned register crossbar can also be called a register cross-switch interconnect, which can be an on-chip direct interconnect path between register files of multiple threads within a thread bundle in a target chip (e.g., a GPU). Through the register crossbar, different threads within a thread bundle can directly access data in the registers of other threads without going through global or shared memory. The register crossbar can be accessed via the thread bundle shuffle instruction, with a latency of only 1 to 2 clock cycles. It does not pass through L1 cache, L2 cache, or global memory, and does not occupy load or store function units, making it an efficient hardware path for data sharing and reduction operations within a thread bundle.

[0084] In some embodiments, after completing the reduction calculation of the current output element, the thread bundle can also iterate through the remaining output elements in an outer loop and execute the corresponding channel calculation tasks sequentially. For example, taking a deformable attention operator, the thread bundle identifier corresponds to a (b, q, h) triple, which includes multiple sampling points (l, p). After completing the reduction calculation of the output element corresponding to the current sampling point, the thread bundle can iterate through the remaining sampling points under (b, q, h) and execute the corresponding channel calculation tasks sequentially.

[0085] Figure 3 This is a schematic diagram of the architecture of a gradient calculation system provided in an embodiment of this disclosure. This gradient calculation system can be used to execute the gradient calculation methods in the foregoing embodiments of this disclosure. Figure 3 As shown, the gradient calculation system may include a conflict topology analysis module 301, a hardware primitive matching module 302, and a thread bundle reduction module 303.

[0086] The conflict topology analysis module 301 can analyze the write conflict pattern of each gradient tensor in the backpropagation gradient tensor list (e.g., candidate gradient tensors to be computed). Based on this write conflict pattern, the gradient tensors can be divided into two categories: the first type can be the target gradient tensor with channel reduction and no write conflict, where each output element of the target gradient tensor is a summation of the channel dimensions, and different output elements are written to different addresses. The second type can be other gradient tensors with spatial scattering and write conflict, where multiple computation sources of the other gradient tensor write to the same target address, resulting in write conflict. This conflict topology analysis module can also be called a conflict topology analyzer.

[0087] For example, for deformable attention operators, the weight gradient and coordinate gradient can belong to the first type of target gradient tensor, and the eigenvalue gradient can belong to other gradient tensors of the second type.

[0088] In some examples, based on this conflict topology analysis module, the write address mapping function of each gradient tensor can be analyzed during the operator development phase to determine the conflict category. For example, for a gradient tensor G, its write address mapping is defined as addr(G, thread_id). If for any two threads t1 != t2 within the same thread bundle, addr(G, t1) = addr(G, t2), and the two threads are responsible for different channel components of the same output element (i.e., reduction is required), then the gradient tensor G is a first-type target gradient tensor. Otherwise, the gradient tensor G is a second-type target gradient tensor.

[0089] For example, the two types can use different criteria for decision-making. For instance, the criteria for decision-making for the first type of target gradient tensor might include: the output dimension of the gradient tensor G does not include the channel dimension c, but the computation depends on the channel dimension data (e.g., it requires reduction of c), and the threads within the same thread bundle have the same target address (corresponding to the same output element). As another example, the criteria for decision-making for other gradient tensors of the second type might include: the output dimension of the gradient tensor G includes the channel dimension c, but multiple input sources map to the same spatial location; thus, threads from different thread bundles might compete for the same address, leading to write address conflicts.

[0090] The hardware primitive matching module 302 can match corresponding hardware primitives for different types of gradient tensors based on the analysis results of the conflict topology analysis module 301. For the first type of target gradient tensor, a reduction method based on the register interconnection network of thread bundles can be selected. For other gradient tensors of the second type, other optimization methods can be used, such as methods based on atomic operations. By selectively matching hardware primitives, the waste of atomic operation resources in write-conflict scenarios can be avoided. This hardware primitive matching module can be implemented based on the kernel. The first type of target gradient tensor and other gradient tensors of the second type can be decomposed into independent kernels, and the optimal hardware primitives are matched for each.

[0091] The thread bundle reduction module 303 can receive the first type of target gradient tensor passed by the hardware primitive matching module 302 and perform reduction operations based on the thread bundle register interconnection network. Based on this thread bundle reduction module, the thread task mapping relationship of the target gradient tensor can be changed: for a target gradient tensor with output elements represented as (b, q, h), it changes from processing one (b, q, h, c) per thread to processing one (b, q, h) per thread bundle. The 32 threads (lanes) within the thread bundle share the calculation of the channel dimension C. For example, thread k processes c = k, k+32, k+64. After each thread accumulates the channel partial sum (thread-level reduction result), the full thread bundle reduction is completed in the register interconnection network through the thread bundle's Shuffle instruction. Only one thread (e.g., thread 0) performs a normal storage write to produce the final result.

[0092] In some examples, Shuffle XOR tree reduction can be used to perform full-thread bundle reduction summation within the register interconnect network. Shuffle XOR tree reduction can merge the channel partial sums of 32 threads into a global sum within the thread bundle register interconnect network (crossbar) through five rounds of __shfl_xor_sync operations (with step sizes of 16, 8, 4, 2, and 1 respectively). The entire process takes approximately 10 clock cycles, generates no memory access requests, does not occupy load / store units, and does not pass through any L1 / L2 / video memory levels.

[0093] The thread bundle reduction module 303 may include a thread bundle granular task allocator, which can reconstruct the thread task mapping from one thread, one element to one thread bundle, one reduction output. In related technologies, the C channel components of the same reduction output are scattered across different thread bundles or even different thread blocks, requiring reduction through global memory atomic operations. In this embodiment, all channel components of the same reduction output can be concentrated in multiple threads within a single thread bundle, and can be directly reduced using the register interconnect network.

[0094] Taking the output element of the target gradient tensor as an example, which includes three dimensions: batch dimension, query position dimension, and attention head dimension, the dimension index value of a certain output element is represented as (b, q, h). The thread mapping rule can be expressed as: Warp ID = (b × Q × H + q × H + h). Here, b is the dimension index value of the batch dimension, q is the dimension index value of the query position dimension, h is the dimension index value of the attention head dimension, Q represents the number of dimension indices for the query position dimension, and H represents the number of dimension indices for the attention head dimension.

[0095] The mapping rules for threads within a thread bundle can include: thread k is responsible for processing channels c such as k, k+W, k+2W, etc., where W is the thread bundle width, i.e. the number of threads contained in the thread bundle.

[0096] Furthermore, for the case where the dimension index value of the output element is represented as (b, q, h, l, p), the outer layer can iterate through the (l, p) sampling points, that is, (l, p) does not participate in the reduction.

[0097] In some examples, the execution process of the thread bundle reduction module 303 may include the following steps S1 to S5.

[0098] Step S1, thread bundle granularity allocation.

[0099] For example, different output elements can be assigned to different thread bundles, such as one thread bundle corresponding to one output element (b, q, h).

[0100] In some examples, this step can convert multiple element representation dimensions of the output element into one-dimensional element indices, and use these one-dimensional element indices as thread bundle identifiers. For instance, for the weight gradient of a deformable attention operator, the batch dimension *b*, query position dimension *q*, and attention head dimension *h* can be converted into one-dimensional element indices, such that one thread bundle corresponds to one (b, q, h) combination. If the batch dimension *B* is 2, the query position dimension *Q* is 10, and the attention head dimension *H* is 8, then 2 × 10 × 8 = 160 thread bundle identifiers can be generated, with each thread bundle responsible for the reduction calculation of one output element.

[0101] Step S2, thread allocation.

[0102] For example, multiple target computation tasks corresponding to the same output element in the target gradient tensor can be assigned to different threads within the same thread bundle for execution.

[0103] In some examples, this step can determine the correspondence between the thread index value of each thread within a thread bundle and the thread allocation index value of the reduced dimension. For instance, if the number of threads in a thread bundle is 32 and the channel dimension C is 32, then the thread numbered k can be responsible for the target computation task with channel index k. If the channel dimension C is 64, then the thread numbered k can be responsible for the target computation tasks with channel indices k and k+32. This allocation method allows the computational load of the channel dimension to be evenly distributed among the threads within the thread bundle.

[0104] Step S3: The results of thread-level reduction are accumulated from the sum of the results of each thread.

[0105] In some examples, each thread in a thread bundle can execute its assigned target computation task based on its corresponding register, calculating at least one thread-level reduction result for a given channel. The thread can then output the thread-level reduction result to a register. For instance, thread number 0 can compute the target computation tasks for channels 0 and 32, summing the results of the two channels and storing the sum in a private register.

[0106] Step S4, reduction operation based on register interconnection network.

[0107] In some examples, this step can perform reduction operations on multiple thread-level reduction results within a thread bundle using a register interconnect network to obtain the element-wise reduction result of the output element. For example, a tree-like reduction can be performed on multiple thread-level reduction results using a thread bundle shuffle instruction. If the thread bundle includes 32 threads, five rounds of XOR shuffle operations can be performed, with step sizes of 16, 8, 4, 2, and 1 for each round. After five rounds, all threads within the thread bundle can hold the same complete reduction result.

[0108] In some examples, this reduction operation can be implemented in hardware, for example, the following five rounds of operations can be implemented in hardware: Round 1: val += __shfl_xor_sync(0xFFFFFFFF, val, 16, 32). Swap thread k with thread k⊕16 and increment (spanning half of the thread bundle).

[0109] Round 2: val += __shfl_xor_sync(0xFFFFFFFF, val, 8, 32). Swap thread k with thread k⊕8 and increment (a quarter of a thread bundle).

[0110] Round 3: val += __shfl_xor_sync(0xFFFFFFFF, val, 4, 32).

[0111] Round 4: val += __shfl_xor_sync(0xFFFFFFFF, val, 2, 32).

[0112] Round 5: val += __shfl_xor_sync(0xFFFFFFFF, val, 1, 32).

[0113] Ultimately, the process is completed through five rounds of reduction, with all threads holding the same global sum. Here, `val` can be the result of the reduction operation on the output elements, and `__shfl_xor_sync` is the XOR shuffle synchronization instruction.

[0114] Step S5: The specified thread performs a normal storage operation.

[0115] In this step, a designated thread within the thread bundle can write the element reduction result to the target chip's memory. For example, the designated thread could be thread number 0 within the thread bundle. This thread can write the complete reduction result to global memory using a non-atomic direct write instruction, while the other threads within the thread bundle do not need to perform the write operation.

[0116] Through the above approach, this gradient computation architecture can match appropriate hardware primitives for different types of gradient tensors based on their write conflict patterns. For target gradient tensors without write conflicts, their reduction operation can be moved from the global storage hierarchy to the intra-thread register interconnect network, with a designated thread writing the reduction result in a non-atomic manner. This eliminates related atomic operations, reduces the processing pressure on the cache controller's CAS queue, and helps improve gradient computation efficiency. Furthermore, this architecture can statically analyze the conflict categories of gradient tensors during the operator development phase, eliminating the need for repeated judgments at runtime and helping to reduce runtime overhead.

[0117] It should be noted that, Figure 3 The gradient calculation system in the illustrated embodiment can achieve Figure 2 The gradient calculation method shown includes all or part of the steps. Figure 2 The specific implementation of the gradient calculation method in the illustrated embodiment can also be found by referring to... Figure 3 The descriptions of the gradient calculation system shown are for reference only, and the embodiments disclosed herein are not limited thereto.

[0118] This disclosure also provides an electronic device, such as... Figure 4As shown, the electronic device 400 may include a memory 410 and a processor 420. The memory can be used to store computer instructions, and the processor can be used to retrieve computer instructions from the memory to execute all or part of the steps of any of the methods in the foregoing embodiments of this disclosure. The processor may be one or more processors, which may execute instructions individually or jointly. Similarly, the memory may be one or more memories, which may store the aforementioned computer instructions individually or jointly. In some examples, the processor may be a CPU, GPU, GPGPU, NPU, or TPU, or it may be a chip in the form of an Application Specific Integrated Circuit (ASIC), System on Chip (SOC), Field Programmable Gate Array (FPGA), Programmable Logic Array (PLA), etc. In some examples, the electronic device may be any type of terminal device, workstation, or server, etc., and this embodiment does not limit this.

[0119] This disclosure also provides a computer-readable storage medium storing a computer program thereon, which, when executed by a processor, implements all or part of the steps of any of the methods in the foregoing embodiments of this disclosure. Optionally, the computer-readable storage medium may be a non-transitory storage medium, but is not limited thereto, it may also be a temporary storage medium.

[0120] This disclosure also provides a computer program product, which may include a computer program that, when executed by a processor, can implement all or part of the steps of any of the methods in the foregoing embodiments of this disclosure.

[0121] The various embodiments in this specification are described in a progressive manner, with each embodiment focusing on the differences from other embodiments. The same or similar parts between the various embodiments can be referred to each other.

[0122] The flowcharts and block diagrams in the accompanying drawings illustrate the architecture, functionality, and operation of possible implementations of systems, methods, and apparatuses according to various embodiments of the present disclosure. In this regard, each block in a flowchart or block diagram may represent a module, unit, or part of a circuit. In some alternative implementations, the functions marked in the blocks may occur in a different order than those marked in the drawings. For example, two consecutive blocks may actually be executed substantially in parallel, or they may sometimes be executed in reverse order, depending on the functions involved. It should be noted that embodiments of the present disclosure may include some or all of the functions marked in the multiple blocks in the drawings, and may also include other functions not shown in the blocks in the drawings. It should also be noted that each block in the block diagrams and / or flowcharts, and combinations of blocks in the block diagrams and / or flowcharts, may be implemented in hardware that performs the specified function or action, or in a combination of dedicated hardware and computer instructions. Unless otherwise specified, implementation in hardware, implementation in software, and implementation in a combination of software and hardware may be equivalent.

[0123] The various embodiments of this disclosure have been described above. These descriptions are exemplary and not exhaustive, and are not limited to the disclosed embodiments. Many modifications and variations will be apparent to those skilled in the art without departing from the scope and spirit of the described embodiments. The terminology used herein is chosen to best explain the principles, practical application, or improvement of the technology in the market, or to enable others skilled in the art to understand the embodiments disclosed herein. The scope of this disclosure is defined by the appended claims.

Claims

1. A gradient calculation method, characterized in that, Applied to a target chip, the target chip includes multiple thread bundles, each thread bundle includes multiple threads, a register corresponding to each thread, and a register interconnection network for interconnection access of multiple registers; The method includes: From the candidate gradient tensors to be calculated, a target gradient tensor that satisfies the preset gradient output conditions is determined; wherein, the preset gradient output conditions include: each of the multiple output elements of the target gradient tensor is generated by a reduction operation based on a set reduction dimension, and different output elements are written to different storage addresses in the memory of the target chip. Multiple target computation tasks corresponding to the same output element in the target gradient tensor are assigned to different threads within the same thread bundle for execution. This allows the thread bundle to execute multiple target computation tasks based on the register and the register interconnection network to perform a reduction operation on the set reduction dimension and obtain the output element. The designated thread within the thread bundle then writes the output element into the memory of the target chip.

2. The method according to claim 1, characterized in that, The output element represents multiple dimensions; the step of assigning multiple target computation tasks corresponding to the same output element in the target gradient tensor to different threads within the same thread bundle includes: Convert the multiple element representation dimensions of the output element into a one-dimensional element index corresponding to the output element; The thread bundle identifier corresponding to the output element is determined based on the one-dimensional element index; For any output element, the multiple target computation tasks corresponding to the output element are assigned to different threads within the thread bundle corresponding to the thread bundle identifier for execution.

3. The method according to claim 2, characterized in that, Converting the multiple element representation dimensions of the output element into a one-dimensional element index corresponding to the output element includes: Determine the dimension index value of the output element in each of the dimensions represented by the element; Obtain the number of dimension indices of the multiple output elements of the target gradient tensor in each of the dimension representations of the elements; Based on the dimension index value and the number of dimension indices, the one-dimensional element index corresponding to the output element is calculated.

4. The method according to claim 1, characterized in that, In a given set of multiple target computation tasks corresponding to the same output element, each target computation task corresponds one-to-one with the reduction dimension index value corresponding to the set reduction dimension of the output element; the step of allocating multiple target computation tasks corresponding to the same output element in the target gradient tensor to different threads within the same thread bundle for execution includes: Determine the number of threads within the thread bundle and the number of reduction dimension indices for the set reduction dimension; Based on the number of threads and the number of reduction dimension indices, determine the correspondence between the thread index value of each thread within the thread bundle and the thread allocation index value of the reduction dimension index; Based on the thread allocation index correspondence, the target computation tasks corresponding to different reduction dimension index values ​​are allocated to the target threads corresponding to the reduction dimension index values.

5. The method according to claim 1, characterized in that, The candidate gradient tensor is the gradient tensor within the target operator. The step of determining the target gradient tensor that satisfies the preset gradient output condition from the candidate gradient tensors to be calculated includes: During the operator development phase, candidate gradient tensors to be computed in the target operator are determined; From the candidate gradient tensors to be calculated, determine the candidate gradient tensors that satisfy the preset gradient output conditions as the target gradient tensor; The method further includes: taking candidate gradient tensors that do not meet the preset gradient output conditions as other gradient tensors; wherein the gradient calculation of the other gradient tensors adopts a different task allocation method than that of the target gradient tensor.

6. The method according to claim 1, characterized in that, The preset gradient output condition also includes: the output dimension of the output element does not include the set reduction dimension.

7. The method according to claim 1, characterized in that, The thread bundle implements the reduction operation for multiple channels in the following ways: Each thread in the thread bundle executes the assigned target computation task based on the register corresponding to the thread to calculate the thread-level reduction result of at least one channel, and outputs the thread-level reduction result to the register; The reduction operation is performed on multiple thread-level reduction results within the thread bundle using a register interconnect network to obtain the element-wise reduction result of the output element.

8. The method according to claim 7, characterized in that, The step of performing reduction operations on multiple thread-level reduction results within the thread bundle through a register interconnect network to obtain the output elements of the target gradient tensor includes: The thread bundle shuffle instruction is executed within the thread bundle through the register interconnection network to perform tree reduction on multiple thread-level reduction results, thereby obtaining the element reduction result of the output element.

9. The method according to claim 1, characterized in that, The target gradient tensor includes the weight gradient and / or coordinate gradient; The reduction dimension includes the channel dimension, and the number of reduction dimension indices of the channel dimension is an integer multiple of the number of threads in the thread bundle.

10. An electronic device, characterized in that, The method includes a memory and a processor, the memory being used to store computer instructions, and the processor being used to retrieve the computer instructions from the memory to perform the method of any one of claims 1 to 9.

11. A computer-readable storage medium, characterized in that, It stores a computer program that, when executed by a processor, implements the method of any one of claims 1 to 9.

12. A computer program product, characterized in that, Includes a computer program that, when executed by a processor, implements the method of any one of claims 1 to 9.

Images