处理单元和矩阵转置复制方法

技术领域

本公开涉及芯片领域，更具体而言，涉及一种处理单元和矩阵转置复制方法。

背景技术

GPU等处理单元中通常包括多个运算单元(例如：核)，各运算单元可分别执行不同的局部运算任务。一些情况下，核在进行局部运算时需要使用存储在处理单元外部的主存储器中的全局数据，例如：在深度神经网络(DNN)反向传播阶段的激活梯度计算任务中，核需要使用存储于主存储器中的源矩阵所对应的转置矩阵作为计算输入条件。

目前，为了将主存储器中的源矩阵所对应的转置后矩阵复制至各运算单元的内存中，常采用的方式为：先通过多级缓存将原始矩阵从主存储器复制到内存，再在各运算单元的内部，通过寄存器将源矩阵从共享内存加载至执行单元，并由执行单元对源矩阵执行转置操作，得到转置后矩阵，最后，再将装置后矩阵返回至内存中。

上述方式，各运算单元在基于转置后矩阵执行后续计算之前，需要先进行转置操作，以获得转置后矩阵。因此，上述方式对处理单元中各运算单元的计算资源消耗较高。

发明内容

有鉴于此，本公开旨在减少矩阵转置复制过程中，对处理单元中各运算单元的计算资源的消耗。

根据本公开的一方面，提供了一种处理单元，包括运算单元、转置单元以及第二存储，其中，

所述运算单元向所述第二存储发送矩阵转置复制指令，所述矩阵转置复制指令包括源地址、目的地地址和转置操作标识；

所述第二存储接收所述矩阵转置复制指令，将与所述源地址中的源矩阵对应的转置后矩阵发送至所述第一存储中的所述目的地地址；

所述转置单元根据所述源地址从所述处理单元外部的主存储器加载所述源矩阵，对所述源矩阵执行转置以得到转置后矩阵。

可选地，所述第二存储在接收到所述矩阵转置复制指令后，确定在所述第二存储中是否已保存所述转置后矩阵；如已保存，将保存的所述转置后矩阵发送到所述目的地地址。

可选地，所述第二存储中存储有所述转置后矩阵以及相应的源地址；所述第二存储通过确定所述第二存储中是否含有与所述源地址相应的转置后矩阵，来确定在所述第二存储中是否已保存所述转置后矩阵。

可选地，所述第二存储通过将所述矩阵转置复制指令中的源地址映射至所述第二存储中的中转地址，并确定所述中转地址处是否存储有数据，以确定在所述第二存储中是否已保存所述转置后矩阵。

可选地，所述第二存储在确定未保存所述转置后矩阵后，将所述矩阵转置复制指令转发到所述转置单元，所述转置单元在对所述源矩阵执行转置以得到转置后矩阵之后，将所述转置后矩阵转发到所述第二存储；其中，所述转置单元将所述源地址映射到所述第二存储中的中转地址，并将所述转置后矩阵转发到所述第二存储中的所述中转地址。

可选地，所述转置单元包括地址映射子单元，用于基于预定映射规则，将所述源地址映射至所述第二存储中的中转地址；其中，所述预定映射规则是固定映射规则或根据所述第二存储中用于存储转置后矩阵的预留存储空间的剩余存储空间大小确定的映射规则；其中，所述映射子单元还用于：在将所述源地址映射至所述第二存储中的中转地址后，确定映射到的中转地址处是否存储数据；如果存储数据，将所述存储的数据移除。

可选地，所述源矩阵为稀疏矩阵，所述稀疏矩阵以数据部分和索引部分存储在所述主存储器中，所述转置单元包括：数据转置子单元，用于结合所述数据部分和所述索引部分，对所述数据部分执行转置操作，得到转置后数据部分；索引转置子单元，用于结合所述数据部分和所述索引部分，对所述索引部分执行转置操作，得到转置后索引部分，其中，所述转置后数据部分和所述转置后索引部分在所述第二存储被合成为所述转置后矩阵。

可选地，所述源矩阵为稠密矩阵，所述稠密矩阵按照预定位置关系划分为多个矩阵块，每个矩阵块具有多行多列，且每个矩阵块中的数据存储在所述主存储器的一个缓存块中，所述转置单元逐缓存块执行转置操作，得到转置后子矩阵，并将所述转置后子矩阵按照预定位置关系合成为转置后矩阵。

根据本公开的第二方面，提供了一种计算装置，包括上述第一方面所述的处理单元。

根据本公开的第三方面，提供了一种矩阵转置复制方法，用于处理单元，所述处理单元包括运算单元、转置单元以及第二存储，所述运算单元具有第一存储，所述方法包括：

通过所述运算单元，向所述第二存储发送矩阵转置复制指令，所述矩阵转置复制指令中包括：源地址、目的地地址和转置操作标识；

通过所述第二存储，将所述矩阵转置复制指令转发到所述转置单元；

通过所述转置单元，根据所述源地址从所述处理单元外部的主存储器加载源矩阵；

通过所述转置单元，对所述源矩阵转置以得到转置后矩阵，将所述转置后矩阵转发到所述第二存储；

通过所述第二存储，将所述转置后矩阵传送到所述第一存储中的所述目的地地址。

可选地，通过所述转置单元，对所述源矩阵转置以得到转置后矩阵，将所述转置后矩阵转发到所述第二存储，包括：

通过所述转置单元，将所述源地址映射到所述第二存储中的中转地址，对所述源矩阵转置以得到转置后矩阵，并将所述转置后矩阵转发到所述第二存储中的所述中转地址。

可选地，所述转置单元包括地址映射子单元；所述通过所述转置单元，将所述源地址映射到所述第二存储中的中转地址，包括：

所述通过所述地址映射子单元，基于预定映射规则，将所述源地址映射至所述第二存储中的中转地址；

通过所述地址映射子单元，确定映射到的中转地址处是否存储数据；如果存储数据，将所述存储的数据移除。

可选地，所述源矩阵为稀疏矩阵，所述稀疏矩阵以数据部分和索引部分存储在所述主存储器中，所述转置单元包括：数据转置子单元和索引转置子单元；

所述通过所述转置单元，对所述源矩阵转置以得到转置后矩阵，将所述转置后矩阵转发到所述第二存储，包括：

通过所述数据转置子单元，结合所述数据部分和所述索引部分，对所述数据部分进行转置，得到转置后数据部分；

通过所述索引转置子单元，结合所述数据部分和所述索引部分，对所述索引部分进行转置，得到转置后索引部分；

将所述转置后数据部分和所述转置后索引部分转发到所述第二存储，其中，所述转置后矩阵在所述第二存储中以所述转置后数据部分和所述转置后索引部分存储。

可选地，所述源矩阵为稠密矩阵，所述稠密矩阵按照预定位置关系划分为多个矩阵块，每个矩阵块具有多行多列，且每个矩阵块中的数据存储在所述主存储器的一个缓存块中；

所述通过所述转置单元，对所述源矩阵转置以得到转置后矩阵，包括：

所述通过所述转置单元，逐缓存块执行转置操作，得到转置后子矩阵，并将所述转置后子矩阵按照预定位置关系合成为转置后矩阵。

本公开实施例中，独立于运算单元之外的转置单元可以根据矩阵转置复制指令中的源地址从处理单元外部的主存储器加载源矩阵，并进行转置操作得到转置后矩阵，以便第二存储将转置后矩阵发送到处理单元的第一存储中，这样，运算单元即可基于第一存储中的转置后矩阵进行后续计算。本公开实施例中，通过转置单元，在对存储于主存储器中的源数据进行复制的过程中进行了转置操作，得到了转置后矩阵，这样，运算单元可以直接基于转置后矩阵执行后续的计算操作，而无需进行转置操作，因此，减少了处理单元中各运算单元的计算资源的消耗。

附图说明

通过参考以下附图对本公开实施例的描述，本公开的上述以及其它目的、特征和优点将更为清楚，在附图中：

图1为本公开一个实施例的处理单元的结构示意图；

图2为稀疏矩阵示意图；

图3为图2所示稀疏矩阵中的数据部分的示意图；

图4为图2所示稀疏矩阵中的索引部分的示意图；

图5为本公开一个实施例的转置单元的结构示意图；

图6为稠密矩阵示意图；

图7为本公开一个实施例的矩阵转置复制方法的流程示意图。

具体实施方式

以下基于实施例对本公开进行描述，但是本公开并不仅仅限于这些实施例。在下文对本公开的细节描述中，详尽描述了一些特定的细节部分。对本领域技术人员来说没有这些细节部分的描述也可以完全理解本公开。为了避免混淆本公开的实质，公知的方法、过程、流程没有详细叙述。另外附图不一定是按比例绘制的。

本公开总体概述

如上所述，GPU等处理单元中包括的各运算单元可用于执行各种不同的计算任务，如深度神经网络中涉及的乘法、加法等一系列运算。具体地：可以从外部主存储器中获取数据，之后再对获取到的数据执行相应地计算任务，从而得到计算结果。

然而在一些计算中，输入数据可能并不是存储于主存储器中的源数据，而是上述源数据对应的转置矩阵，例如：执行深度神经网络反向传播阶段的激活梯度计算任务时，计算的输入数据是存储于主存储器中的源矩阵所对应的转置矩阵，而并非上述源数据本身。在上述情况下，则需要对存储于主存储器中的源矩阵执行转置操作得到转置矩阵，并将转置矩阵复制至运算单元内部，以供后续计算使用。

为了在运算单元内部得到转置矩阵，常采用的方式为，先通过多级缓存将源矩阵从主存储器复制到运算单元内存，再在运算单元的内部，通过寄存器将源矩阵从内存加载至执行单元，并由执行单元对源矩阵执行转置操作，得到转置后矩阵，最后，再将装置后矩阵返回至内存中。上述方式，运算单元在基于转置后矩阵执行后续计算之前，需要先进行转置操作，以获得转置后矩阵。因此，上述方式对运算单元内部的计算资源消耗较高。

而本公开实施例中，处理单元中独立于运算单元之外的转置单元可以对源矩阵进行转置操作得到转置后矩阵，以将转置后矩阵转发至运算单元内部，这样，运算单元即可基于内部存储的转置后矩阵进行后续计算，而无需进行转置操作，因此，减少了对处理单元中各运算单元计算资源的消耗。

本公开实施例的详细实现过程

参见图1，图1为本公开一个实施例的处理单元的结构示意图。其中，处理单元100包括：运算单元110、转置单元120以及第二存储130，所述运算单元110具有第一存储1101。运算单元110向所述第二存储130发送矩阵转置复制指令，矩阵转置复制指令包括源地址、目的地地址和转置操作标识。第二存储130接收所述矩阵转置复制指令，将与源地址中的源矩阵对应的转置后矩阵发送至所述第一存储1101中的目的地地址。转置单元120根据源地址从处理单元100外部的主存储器200加载源矩阵，对源矩阵执行转置以得到转置后矩阵。

本公开实施例中，处理单元100指进行图像处理以及图形相关运算处理的单元，如进行深度神经网络中激活梯度计算处理的GPU(graphics processing unit，图形处理器)等。运算单元110为用于执行具体的运算任务并输出运算结构的单元。

处理单元100内部的运算单元110的个数可以是一个也可以是多个(参见图1所示)，当运算单元110的数量为多个时，各运算单元110可以分别执行不同的运算任务，本公开中对于运算单元110的数量不做限定。

第二存储130可以为位于各运算单元110之间的存储单元，例如，可以为缓存等；第一存储1101可以为运算单元110内部所具有的存储单元，例如：也可以为缓存等。

矩阵转置复制指令是用于将主存储器200源地址处的源矩阵转置，以便复制到第一存储1101的目的地地址中的指令。矩阵转置复制指令中可以至少包括：源地址、目的地地址和转置操作标识，另外，也可以包括复制操作标识。本公开实施例中，对于矩阵转置复制指令的具体格式不做限定，示例性地，矩阵转置复制指令的格式可以为：“cp.async.trans[dst][src]”其中，“src”表示源地址；“dst”表示目的地地址；“trans”表示转置操作标识；“cp.async”为复制操作标识。

具体地，矩阵转置复制指令中的源地址，可以表示在处理单元100外部的主存储器200中的源矩阵的存储地址，其中，源矩阵为转置单元120所针对的待进行转置操作的矩阵，目的地地址则可以表示在运算单元110的第一存储1101中要存储的转置后矩阵的存储地址。

就矩阵而言，按照其中包括的0元素与非0元素的比例的不同，可以分为稀疏矩阵和稠密矩阵，其中，稀疏矩阵是指在矩阵中，数值为0的元素数目远远多于非0元素的数目，并且非0元素分布没有规律的矩阵。与之相反，稠密矩阵则是指非0元素数目占大多数的矩阵。本公开实施例中的源矩阵可以为稀疏矩阵，也可以为稠密矩阵。

就处理单元100中的运算单元110而言，在使用转置后矩阵进行某具体运算之前，需要先获取到上述转置后矩阵，亦即，需要将上述转置后矩阵存储在内部存储单元第一存储1101中。

具体地，处理单元100中的运算单元110可以向第二存储130发送矩阵转置复制指令，该矩阵转置复制指令中包括：源地址、目的地地址和转置操作标识。第二存储130在接收到运算单元110发送的上述矩阵转置复制指令之后，则可以在获取到源矩阵对应的转置后矩阵之后，将转置后矩阵发送至第一存储1101中的目的地地址，以便运算单元110使用该转置后矩阵进行后续的具体运算。具体地，本公开实施例中，对于后续具体运算的具体形式不做限定，例如，可以对转置后矩阵与其他矩阵进行矩阵的乘法运算，也可以对转置后矩阵与其他矩阵进行矩阵的加法运算，还可以先将转置后矩阵与某个矩阵(例如矩阵A)进行矩阵的乘法运算，得到乘法运算结果，再将乘法运算结果与另一个矩阵(例如矩阵B)进行矩阵的加法运算，等等。处理单元100中的转置单元120可以在接收到矩阵转置复制指令之后，根据其中包括的源地址从主存储器200加载源矩阵，对源矩阵执行转置以得到转置后矩阵。

本公开实施例中，独立于运算单元110之外的转置单元120可以根据矩阵转置复制指令中的源地址，从处理单元100外部的主存储器200加载源矩阵，并进行转置操作得到转置后矩阵，以便第二存储130将转置后矩阵发送到处理单元100的第一存储1101中，运算单元110即可基于第一存储1101中的转置后矩阵进行后续计算。本公开实施例中，通过转置单元120，在对存储于主存储器200中的源数据进行复制的过程中进行了转置操作，得到了转置后矩阵，运算单元110可以直接基于转置后矩阵执行后续的计算操作，而无需进行转置操作，因此减少了处理单元100中各运算单元110的计算资源的消耗。

可选地，在本公开的一些实施例中，第二存储130在接收所述矩阵转置复制指令之后，可以先确定在本地是否已保存有源矩阵对应的转置后矩阵，如已保存，则可以将保存的转置后矩阵发送到目的地地址。

第二存储130的作用是接收矩阵转置复制指令，并将与源地址中的源矩阵对应的转置后矩阵发送至第一存储中的所述目的地地址。因此，当第二存储130接收了较多的矩阵转置复制指令后，其中可能已经存储有多个矩阵转置复制指令的源矩阵对应的转置后矩阵，上述多个转置后矩阵可以为转置单元对源矩阵执行转置得到转置后矩阵之后转发至第二存储130的。此种情况下，当第二存储130接收到新的矩阵转置复制指令后，其本地可能已经存储有该新的矩阵转置复制指令的源矩阵对应的装置后矩阵，此时，可以先确定本地是否保存有转置后矩阵。如此一来，则可以避免转置单元120对源矩阵再次执行转置操作，节省了运算资源。同时，由于无需执行转置操作，而是直接将第二存储130存储的转置后矩阵发送至第一存储1101中，因此获取转置后矩阵的效率较高。

具体地，第二存储130确定是否已保存源矩阵对应的转置后矩阵的方式可以为：在第一种方式，第二存储可以在存储转置后矩阵的同时，还存储相应的源矩阵的源地址，进而通过存储的源地址来确定是否保存有源矩阵对应的转置后矩阵。具体地，第二存储130可以通过确定所述第二存储130中是否含有与所述源地址相应的转置后矩阵，来确定在所述第二存储130中是否已保存所述转置后矩阵。

在第二种方式，第二存储130在接收到矩阵转置复制指令之后，再将源地址映射至第二存储130的中转地址，再根据中转地址中是否有数据，确定是否保存有转置后矩阵，其中，中转地址为第二存储130中预先保留的用于存储源矩阵对应的转置后矩阵的存储区域的地址。具体地，第二存储130通过将矩阵转置复制指令中的源地址映射至第二存储130中的中转地址，并确定中转地址处是否存储有数据，以确定在第二存储130中是否已保存转置后矩阵。

上述两种方式相比，第一种方式中预选存储了转置后矩阵以及相应的源地址，这样，第二存储130在接收到矩阵转置复制指令后，可以直接根据源地址确定出需要的转置后矩阵，而无需进行其他操作，因此效率较高。而第二种方式中，第二存储中仅需存储转置后矩阵，而无需存储源地址，因此，与第一种方式相比，减少了对存储空间的占用，可以节省存储资源。

可选地，在本公开的一些实施例中，当第二存储130确定未保存转置后矩阵，则可以将矩阵转置复制指令转发到所述转置单元120，转置单元120在对源矩阵执行转置以得到转置后矩阵之后，再将转置后矩阵转发到第二存储130，以便第二存储130将与转置后矩阵发送至所述第一存储1101中的目的地地址。

进一步地，转置单元120可以将接收到的矩阵转置复制指令中的源地址映射到所述第二存储130中的中转地址，之后，再将转置后矩阵转发到第二存储130中的中转地址。具体地，转置单元120中可以包括地址映射子单元，该地址映射子单元即用于基于预定映射规则，将源地址映射至所述第二存储130中的中转地址。

本公开实施例中，对于映射规则的具体内容不做限定，可以根据实际情况进行设定。具体地，可以是按照某种固定的映射规则映射得到中转地址，例如，可以按照第二存储130接收到矩阵转置复制指令的先后顺序，依次在第二存储130中确定各矩阵转置复制指令中源地址对应的中转地址。亦即前一个矩阵转置复制指令中的源地址对应的中转地址，与后一个矩阵转置复制指令中的源地址对应的中转地址，是两个相邻的地址。也可以根据源地址在主存储器中的位置，确定第二存储130中中转地址的位置等等。另外，映射规则也可以是根据第二存储130中用于存储转置后矩阵的预留存储空间的剩余存储空间大小确定的映射规则。

在转置单元120中的映射子单元将源地址映射至所述第二存储130中的中转地址后，中转地址中可能还存储有其他的数据，为避免因中转地址中存在其他数据而导致转置单元120得到的转置后矩阵无法存储问题的出现，在其中一些实施例中，映射子单元在将源地址映射至所述第二存储130中的中转地址后，还可以先确定映射到的中转地址处是否存储有数据。如果未存储有数据，则转置单元120可以在对源矩阵执行转置以得到转置后矩阵之后，将转置后矩阵直接存储至第二存储130的中转地址中。如果中转地址中已经存储有数据，则映射子单元可以将已存储的数据移除，以便转置单元120得到转置后矩阵之后，将转置后矩阵存储至第二存储130的中转地址中。

可选地，在其中一些实施例中，源矩阵为稀疏矩阵，并且所述稀疏矩阵是以数据部分和索引部分存储在主存储器中的，其中，数据部分即为稀疏矩阵中的非0元素部分，索引部分则是指用于表示各非0元素在整个稀疏矩阵中的位置信息的部分。例如在图2-图4中，图2所示即为稀疏矩阵，图3所示即为图2所述的稀疏矩阵中的数据部分，其中，包括非0元素：A、B、C、D、E、F、G以及H，图4所示即为图2所述的稀疏矩阵中的索引部分，用于表示图3所示的数据部分在图2所示的整个稀疏矩阵中的位置，例如：与A对应的0表示：A位于第0行第0列，与B对应的1表示：B位于第0行第1列，……，依次类推。

参见图5所示，当源矩阵为稀疏矩阵时，转置单元可以包括：数据转置子单元1201和索引转置子单元1202，其中，数据转置子单元1201用于结合数据部分和索引部分，对数据部分执行转置操作，得到转置后数据部分。索引转置子单元1202，用于结合数据部分和索引部分，对索引部分执行转置操作，得到转置后索引部分。其中，转置后数据部分和转置后索引部分在第二存储130被合成为转置后矩阵。

具体地，例如可以结合数据部分和索引部分，先将以数据部分和索引部分存储的稀疏矩阵还原为以图4所示的常规形式的稀疏矩阵，然后针对还原后稀疏矩阵，沿对角线将对称分布的各元素进行交换，得到交换后矩阵。之后，再由数据转置子单元1201，基于交换后矩阵得到转置后数据部分，由索引转置子单元1202基于交换后矩阵得到转置后索引部分。

可选地，在其中一些实施例中，源矩阵为稠密矩阵，所述稠密矩阵按照预定位置关系划分为多个矩阵块，每个矩阵块具有多行多列，且每个矩阵块中的数据存储在主存储器200的一个缓存块中，转置单元120逐缓存块执行转置操作，得到转置后子矩阵，并将转置后子矩阵按照预定位置关系合成为转置后矩阵。

具体地，参见图6所示，图6左侧为一个稠密矩阵，右侧为左侧稠密矩阵对应的转置后矩阵。存储时，由于通常存储器中每个缓存块存储的数据量较为有限，因此一个完整的矩阵可以被划分为多个缓存块进行存储的，每个缓存块存储一定数量的元素。现有技术中，通常是按照矩阵中行的顺序进行存储的，例如左侧的稠密矩阵，第一个缓存块存储第一行的元素：A、B、C、D；第二个缓存块存储第二行的元素：E、F、G、H；第三个缓存块存储第三行的元素：I、J、K、L；第四个缓存块存储第四行的元素：M、N、O、P。这样一来，在转置单元120在对源矩阵进行转置的过程中，以得到右侧的第一行元素为例，则需要分别从第一个缓存块读取元素A、从第二个缓存块读取元素E、从第三个缓存块读取元素I、从第四个缓存块读取元素M，也就是说，为了得到第一行的所有元素，转置单元120需要跨4个缓存块进行数据读取。

而本公开实施例中，则将稠密矩阵按照预定位置关系划分为多个矩阵块，每个矩阵块具有多行多列，且每个矩阵块中的数据存储在所述主存储器200的一个缓存块中。仍以图6进行说明，本公开中，例如，可以将A、B、E以及F组成的第一矩阵块存储在同一缓存块(第一缓存块)中，将C、D、G以及H组成的第二矩阵块存储在同一缓存块(第二缓存块)中，将I、J、M以及N组成的第三矩阵块存储在同一缓存块(第三缓存块)中，将K、L、O以及P组成的第四矩阵块存储在同一缓存块(第四缓存块)中。这样转置单元120可以逐缓存块分别对每个缓存块中存储的矩阵块进行转置，得到转置后子矩阵，进而合并成转置后矩阵。例如，针对第一矩阵块，可以仅读取第一缓存块，而无需跨线程即可得到第一矩阵块，之后，对该第一矩阵块进行转置操作，得到转置后子矩阵。

相比可知，本公开实施例中，转置单元120在执行转置操作的过程中，仅需读取一个缓存块，而无需跨缓存块，即可得到一个完整的矩阵块，因此，本公开实施例有效提升了转置操作的效率。

可选地，在其中一些实施例中，稀疏矩阵的数据部分和索引部分存储在主存储器200的一个缓存块中，数据转置子单元1201和索引转置子单元1201基于该缓存块中的数据部分和索引部分，分别得到转置后数据部分和转置后索引部分。

针对稀疏矩阵而言，无论是得到转置后数据部分，还是得到转置后索引部分，都需要同时获取到稀疏矩阵的数据部分和索引部分。因此，本公开实施例中，将数据部分和索引部分存储在主存储器200的一个缓存块中，这样，数据转置子单元1201可以从同一个缓存块中同时获取数据部分和索引部分，而无需跨缓存块读取。同样地，索引转置子单元1202也可以从同一个缓存块中同时获取数据部分和索引部分，而无需跨缓存块读取，因此本公开实施例有效提升了转置操作的效率。

如图7所述，本公开的一个实施例还提供了一种矩阵转置复制方法，用于处理单元100，所述处理单元100包括运算单元110、转置单元120以及第二存储130，运算单元110具有第一存储1101，所述方法包括：

步骤701，通过运算单元110，向第二存储130发送矩阵转置复制指令，矩阵转置复制指令中包括：源地址、目的地地址和转置操作标识。

步骤702，通过第二存储130，将矩阵转置复制指令转发到转置单元120。

步骤703，通过转置单元120，根据源地址从处理单元100外部的主存储器200加载源矩阵。

步骤704，通过转置单元120，对源矩阵转置以得到转置后矩阵，将转置后矩阵转发到第二存储130。

步骤705，通过第二存储130，将转置后矩阵传送到第一存储1101中的目的地地址。

可选地，通过转置单元120，对源矩阵转置以得到转置后矩阵，将转置后矩阵转发到第二存储130，包括：通过转置单元120，将源地址映射到第二存储130中的中转地址，对源矩阵转置以得到转置后矩阵，并将转置后矩阵转发到第二存储130中的中转地址。

可选地，转置单元120包括地址映射子单元；通过所述转置单元120，将源地址映射到第二存储130中的中转地址，包括：通过地址映射子单元，基于预定映射规则，将源地址映射至第二存储130中的中转地址。

可选地，在通过地址映射子单元，基于预定映射规则，将源地址映射至第二存储130中的中转地址后，所述方法还包括：通过地址映射子单元，确定映射到的中转地址处是否存储数据；如果存储数据，将存储的数据移除。

可选地，所述源矩阵为稀疏矩阵，稀疏矩阵以数据部分和索引部分存储在主存储器200中，转置单元120包括：数据转置子单元和索引转置子单元。在此实施例中，通过所述转置单元120，对源矩阵转置以得到转置后矩阵，将转置后矩阵转发到所述第二存储130，包括：通过数据转置子单元，结合数据部分和索引部分，对数据部分进行转置，得到转置后数据部分。通过索引转置子单元，结合数据部分和索引部分，对索引部分进行转置，得到转置后索引部分。将转置后数据部分和转置后索引部分转发到第二存储130，其中，转置后矩阵在第二存储130中以转置后数据部分和转置后索引部分存储。

可选地，所述源矩阵为稠密矩阵，稠密矩阵按照预定位置关系划分为多个矩阵块，每个矩阵块具有多行多列，且每个矩阵块中的数据存储在所述主存储器200的一个缓存块中。在此实施例中，通过转置单元120，对源矩阵转置以得到转置后矩阵，包括：通过转置单元120，逐缓存块执行转置操作，得到转置后子矩阵，并将转置后子矩阵按照预定位置关系合成为转置后矩阵。

可选地，稀疏矩阵的数据部分和索引部分存储在主存储器200的一个缓存块中。在此实施例中，通过所述数据转置子单元，结合数据部分和索引部分，对数据部分进行转置，得到转置后数据部分，包括：通过数据转置子单元，基于该缓存块中的数据部分和索引部分，对数据部分进行转置，得到转置后数据部分。通过索引转置子单元，结合数据部分和索引部分，对索引部分进行转置，得到转置后索引部分，包括：通过索引转置子单元，基于该缓存块中的数据部分和索引部分，对索引部分进行转置，得到转置后索引部分。

由于上述过程的实现细节已在前述装置实施例的描述中详细介绍，故不赘述。

应该理解，本说明书中的各个实施例均采用递进的方式描述，各个实施例之间相同或相似的部分互相参见即可，每个实施例重点说明的都是与其他实施例的不同之处。尤其，对于方法实施例而言，由于其基本相似于装置和系统实施例中描述的方法，所以描述的比较简单，相关之处参见其他实施例的部分说明即可。

应该理解，上述对本说明书特定实施例进行了描述。其它实施例在权利要求书的范围内。在一些情况下，在权利要求书中记载的动作或步骤可以按照不同于实施例中的顺序来执行并且仍然可以实现期望的结果。另外，在附图中描绘的过程不一定要求示出的特定顺序或者连续顺序才能实现期望的结果。在某些实施方式中，多任务处理和并行处理也是可以的或者可能是有利的。

应该理解，本文用单数形式描述或者在附图中仅显示一个的元件并不代表将该元件的数量限于一个。此外，本文中被描述或示出为分开的模块或元件可被组合为单个模块或元件，且本文中被描述或示出为单个的模块或元件可被拆分为多个模块或元件。

还应理解，本文采用的术语和表述方式只是用于描述，本说明书的一个或多个实施例并不应局限于这些术语和表述。使用这些术语和表述并不意味着排除任何示意和描述(或其中部分)的等效特征，应认识到可能存在的各种修改也应包含在权利要求范围内。其他修改、变化和替换也可能存在。相应的，权利要求应视为覆盖所有这些等效物。