CN117785287A - 多线程计算中的私有存储器模式顺序存储器访问 - Google Patents

Abstract

提供了用于线程级并行处理(其中warp中的线程被并行执行)的处理器、系统和方法。一种方法可以包括：序列发生器向存储器端口发送序列发生器的标量寄存器中存储的线程块大小信息和寻址参数，存储器端口基于线程块大小信息和寻址参数生成标量存储器地址，存储器接口和存储器单元从起始于标量存储器地址的K个连续存储器地址将warp的数据的K个字加载到向量寄存器。标量存储器地址可以指向warp的第一线程的数据的存储位置。并且K可以是warp大小，warp大小表示要通过执行向量指令来同时处理的线程数。

Description

多线程计算中的私有存储器模式顺序存储器访问

技术领域

本公开涉及计算机架构，尤其涉及一种在单指令多线程(single instructionmultiple threads，SIMT)计算系统中使用标量存储器地址来加载和存储多线程的向量数据的多线程计算机架构。

背景技术

图形处理单元(graphics processing unit，GPU)架构提供了一种以单指令多线程(single instruction multiple threads，SIMT)方式执行并行线程的方法。诸如GPU的SIMT处理器具有多个用于同时执行多个线程的核，并且特别适合于大规模并行计算应用。为利用上述用于并行执行的多个核，计算机程序通常需要通过调用专门为多核工作而设计的应用编程接口(application programming interface，API)的函数来为多核架构进行定制。最近，利用GPU在传统上由中央处理器(central processing unit，CPU)处理的应用中执行计算的通用计算GPU(general-purpose computing on GPU，GPGPU)变得更加实用和流行。然而，为多个并发线程加载和存储数据需要每个线程的数据地址。上述数据地址通常需要事先通过向量运算来准备，并且需要使用向量地址总线。因此，本领域需要更有效地为多个线程加载和存储数据的处理器。

发明内容

本公开描述了使用标量存储器地址有效地加载向量数据的装置、方法和系统。在一示例性实施例中，一种方法可以包括：序列发生器向存储器端口发送序列发生器的标量寄存器中存储的线程块大小信息和寻址参数，存储器端口基于线程块大小信息和寻址参数生成标量存储器地址，以及存储器接口和存储器单元从起始于标量存储器地址的K个连续存储器地址将warp的数据的K个字加载到向量寄存器。标量存储器地址可以指向数据warp的第一线程的数据的存储位置。K可以是warp大小，warp大小表示要通过执行向量指令来同时处理的线程数。该数据的大小可以小于或等于字大小。

在另一示例性实施例中，处理器可以包括序列发生器、被耦合至该序列发生器的存储器端口、以及经由存储器接口耦合至该存储器端口的存储器单元。序列发生器可以用于向存储器端口发送序列发生器的标量寄存器中存储的线程块大小信息和寻址参数，并且存储器端口可以用于基于线程块大小信息和寻址参数生成标量存储器地址，所述标量存储器地址指向数据warp的第一线程的数据的存储位置。存储器单元和存储器接口可以用于：对于数据加载过程，从起始于标量存储器地址的K个连续存储器地址将warp的数据的K个字加载到向量寄存器，或者对于数据存储过程，从向量寄存器将warp的数据的K个字存储到起始于标量存储器地址的K个连续存储器地址。K可以是warp大小，warp大小表示要通过执行向量指令来同时处理的线程数。该数据的大小可以小于或等于字大小。

在另一实施例中，一种方法可以包括：序列发生器向存储器端口发送序列发生器的标量寄存器中存储的线程块大小信息和寻址参数；存储器端口基于线程块大小信息和寻址参数生成标量存储器地址；存储器接口和存储器单元从向量寄存器将warp的数据的K个字存储到起始于所述标量存储器地址的K个连续存储器地址。标量存储器地址可以指向数据warp的第一线程的数据的存储位置，并且K可以是warp大小，warp大小表示要通过执行向量指令来同时处理的线程数。该数据的大小可以小于或等于字大小。

附图说明

图1示意性地示出了根据本公开实施例的处理器。

图2示意性地示出了根据本公开实施例的处理器的处理元件。

图3示意性地示出了根据本公开实施例的处理器的存储器单元。

图4A示意性地示出了根据本公开实施例的存储器接口。

图4B示意性地示出了根据本公开实施例的私有存储器访问接口。

图4C示意性地示出了根据本公开实施例的共享存储器访问接口。

图5示意性地示出了根据本公开实施例的处理器的交换盒(switch box)。

图6示意性地示出了根据本公开实施例的处理器的垫片存储器。

图7示意性地示出了根据本公开实施例的处理器的存储器端口。

图8示意性地示出了根据本公开实施例的处理器的序列发生器。

图9是根据本公开实施例的加载或存储的多线程内核的数据的方法的流程图。

具体实施方式

现在将详细地介绍本教导的各实施例，这些实施例的示例在附图中示出。为了一致性，不同附图中的相似元素用相似的参考标号表示。尽管将结合各实施例来描述本教导，但可以理解，本教导不仅限制于这些实施例。相反，本教导旨在涵盖在所附权利要求中定义的精神和范围内的替代品、修改、以及等同物。

此外，在本公开的实施例的下面的具体实施方式中，阐述了许多具体细节，以提供对本教导的彻底理解。然而，本领域的普通技术人员会意识到，没有这些具体细节，本教导也可以被实施。在其他情况下，没有详细描述众所周知的方法、程序、组件、以及电路，以免不必要地模糊了本教导的实施例的各个方面。

图1示意性地示出了根据本公开的实施例的处理器100。处理器100可以包括直接存储器访问(direct memory access，DMA)模块102、配置存储器104、序列发生器106、控制处理器108、程序高速缓存110、存储器单元112、多个存储器接口124.1-124.N、PE阵列114和垫片存储器116。DMA模块102可以耦合到外部总线130，并且可以由控制处理器108控制。DMA模块102可以负责将可执行指令和不可执行数据移入和移出外部总线130。程序高速缓存110可以存储由控制处理器108用来控制DMA模块102的操作的指令和数据。在一个实施例中，控制处理器108可以使用存储在程序高速缓存110中的指令和数据来处理内核程序(kernel program)。

在一些实施例中，处理器100可以被配置用于大规模线程级并行处理。例如，PE阵列114中的一个处理元件(processing element，PE)可以包括向量大小大于1的向量算术逻辑单元(arithmetic logic unit，ALU)，并且向量ALU的每个ALU可以用于对不同的数据执行相同的操作(例如，每个线程可以对其自己的数据进行操作)。也就是说，在具有多个ALU的这些实施例中，每个PE可以被配置成以单指令多线程(single instruction multiplethreads，SIMT)方式操作。在一个实施例中，PE可以通过多个向量数据输入生成一个向量数据输出。在一些实施例中，线程也可以称为流。

为了给要被同时执行的多个线程提供数据，在一些实施例中，处理器100的组件之间的一些相关电子连接可以是向量型式。例如，KxG位的向量地址连接(vector addressconnection)可以具有K数量的G位地址，而KxM位的向量数据连接可以具有K数量的M位数据。还应注意，尽管在任何附图中未示出，但不同组件之间的数据或地址连接可伴随有一个或多个控制信号线。例如，繁忙信号线可以存在于第一组件和第二组件之间，并且可以被第一组件用来向第二组件发送繁忙信号，指示第一组件没有准备好接受有效的数据或地址信号。此外，有效信号线也可以存在于第一组件和第二组件之间，并且可以被第二组件用来向第一组件发送有效信号，指示有效数据或地址信号已经被放置在连接线上。

配置存储器104可以存储包括算术和逻辑指令的数据路径程序(data pathprogram)，并且加载和存储用于数据路径组件的指令。在一个实施例中，存储在配置存储器104中的数据路径程序可以是已编译指令的一个或多个序列。例如，数据路径程序可以包括要由PE阵列114执行的指令，该指令可以指定PE可以执行什么种类的操作，以及数据路径组件可以如何保持数据或发送数据。

序列发生器106可以对存储在配置存储器104中的指令进行解码。指令可以包括标量指令和向量指令。对于标量指令，序列发生器106可以解码标量指令并执行由标量指令编码后的标量操作。对于向量指令，序列发生器106可以对向量指令进行解码，并将解码后的向量指令传递给PE阵列114的各种组件(例如，PE阵列114的将涉及算术和逻辑操作以及数据移动操作的组件)，使得由向量指令编码的向量操作可以由PE阵列114的组件执行。PE阵列114的这些组件也可以被称为向量处理单元。如本文所使用的，标量操作可以指标量指令的执行，而向量操作可以指向量指令的执行。

解码后的向量指令可以被打包传递到各种组件，这些打包的向量指令可以被称为配置包或简称为配置。除了解码后的指令之外，一个组件的配置包可以包括一些其他参数(例如，warp的数量(指定在一个配置设置中，指令要被重复执行多少次以及数据通过数据交换单元多少次)、以及从序列发生器传递到组件的立即数(immediate value))。如本文所使用的，一个warp可以指一个PE中并发执行的线程数量，例如，对于有32个ALU的PE，一个warp可以指由PE执行的32个线程。在一个实施例中，物理数据路径配置可以被称为物理数据路径程序，其可以包括物理数据路径中包括的各种组件的单独配置。

尽管未示出，但可以存在将序列发生器106连接到PE阵列114的组件的配置总线，可以用于经由总线将单独配置传递给这些组件。例如，存在可以用于向存储器端口、交换盒、以及PE传递配置的配置总线。在一些实施例中，存储器端口的配置可以包括数据准备指令，例如但不限于加载/存储指令(以及指令的参数，例如地址)，并且PE的配置可以包括要由PE中的ALU执行的指令，例如但不限于像加法或减法的数据处理指令。

存储器单元112可以是数据暂存区域(data staging area)，用于存储从外部总线130接收的数据以及由PE阵列114生成的执行结果数据(在这些结果可能经由外部总线130发送出去之前)。在一些实施例中，存储器单元112可以是处理器100外部的大型存储器系统的处理器内高速缓存。

PE阵列114可以包括多个存储器端口(memory port，MP)120.1-120.N、多个交换盒(switch box，SB)122.1-122.N、以及多个处理元件(processing element，PE)118.1-118.N。这些组件可形成N列的可编程硬件单元或可编程硬件组件。例如，MP 120.1、SB122.1和PE118.1可以形成PE阵列114的第一列，并且MP 120.N、SB 122.N和PE 118.N可以形成PE阵列114的第N列。在这些硬件单元中的每一个被配置用于向量处理的实施例中，可编程硬件单元的每个列也可以被称为向量处理单元的一列或简称为一列。在图1所示的示例中，PE阵列114可以包括一行处理元件118.1-118.N，其可以包括N个PE，其中N是整数。在本文描述的实施例中，数字N可以是32。但这是示例，在其他实施例中数字N可以是另一整数，例如但不限于16、64等。尽管给出的这些示例数字可以是2的幂，但在一个实施例中，一个行中的PE的数量不必是2的幂。

多个MP 120.1-120.N可以是控制PE阵列114和存储器单元112之间的数据流的可编程硬件单元。如图1所示，多个MP 120.1-120.N可以经由多个存储器接口(memoryinterface，MI)124.1-124.N耦合到存储器单元112。每个MP 120.1-120.N可以经由相应的MI124.1-124.N耦合到存储器单元112，以从存储器单元112读取和写入存储器单元112。因此，MI 124.1可以是到PE阵列114的第一列的存储器单元112的网关，以此类推，直到MI124.N可以是到PE阵列114的第N列的存储器单元112的网关。每个MP 120.1-120.N还可以耦合到相应列中的SB，以向每个列发送数据和从每个列发送数据。例如，MP 120.1可以耦合到SB 122.1，MP 120.2可以耦合到SB 122.2，等等。在一些实施例中，存储器单元112和MI124.1-124.N可以统称为多端口存储器单元。此外，在至少一个实施例中，存储器单元112、MI 124.1-124.N和MP 120.1-120.N可以支持两种存储器访问模式：私有存储器访问模式和共享存储器访问模式，其也可以被称为私有存储器访问方法(或简称为私有存储器访问)和共享存储器访问方法(或简称为共享存储器访问)。应当注意，虽然MI 124.1-124.N在图1中被示为单独的实体，但是在一些实施例中，MI 124.1-124.N可以分别被实现为相应MP120.1-120.N的一部分，或者在一些其他实施例中，被实现为存储器单元112的一部分。

除了MP 120.1和MP 120.N之外，所有MP都可以耦合到两个相邻MP，使得每个MP可以被配置成从前续MP接收地址和/或向后续MP发送地址。MP之间的电子耦合可以提供地址的单向流(例如，在一个配置中，指定地址可以从一个MP流向下一MP)。例如，如图1所示，MP120.1可以耦合到MP 120.2用于单向地址流，MP 120.2可以耦合到MP 120.3用于单向地址流。最后一个MP 120.N可以是例外，并且耦合到垫片存储器116，该垫片存储器116可以为地址提供临时存储。第一个MP 120.1可以是另一个例外，因为第一个MP 120.1可以从垫片存储器116接收单向地址流。在一些实施例中，MP 120.1-120.N可以沿着PE行方向形成一个地址路由总线。即，地址可以在MP之间以与数据可以以其在PE和SB之间路由的方向的平行方向上路由。在至少一个实施例中，在MP之间传输的地址可以是由MP更新的存储器地址。例如，MP可以对存储器地址执行一些操作(例如，更新存储器地址)并将更新的存储器地址传送给后续列中的下一MP。

垫片存储器116可以用作数据缓冲器，例如先进先出(first-in-first-out，FIFO)，以从PE阵列(例如，从MP 120.N、PE 118.N和/或SB 122.N)收集地址和数据，并且当PE阵列的第一列(例如，MP 120.1、SB 122.1和PE 118.1)被新配置重配置时，将这些地址和数据反馈回PE阵列(例如，到MP 120.1和/或SB 122.1)。

在一些实施例中，可以在处理线程块时静态地配置PE、SB和MP。例如，PE、SB和MP中的每一个可以用在相应配置中指定的指令来编程，以执行流水线(pipeline)的一个级(stage，也称为阶段)。当线程块中的数据正通过流水线级时，不可以更改任何指令。例如LOAD和STORE之类的地址计算指令和存储器访问指令可以被映射到MP(例如，封装在发送到相应MP的配置中)，数据交换指令可以被映射到SB(例如，封装在发送到相应SB的配置中)，并且其他指令可以被映射到PE(例如，封装在发送到相应PE的配置中)。

如图1所示，除了SB 122.1和SB 122.N之外，SB 122.1-122.N中的每一个都可以耦合到两个相邻的SB(例如，前续SB和后续SB)。SB 122.1可以耦合到MP 120.1、垫片存储器116、PE 118.1和SB 122.2。并且SB 122.N可以耦合到MP 120.N、垫片存储器116、PE 118.N和SB 122.N-1。在一些实施例中，SB 122.1-122.N可以沿着PE行方向形成数据路由总线。即，数据可以在SB之间以与数据可以以其在PE之间路由的方向的平行方向上路由。在一个实施例中，一个或多个SB可以用于路由数据以绕过一个或多个PE。

SB 122.1-122.N可以用于提供用于以下数据的数据交换：在相邻PE之间路由的数据、从PE路由到MP的数据，从PE路由到数据路由总线的数据、从MP路由到PE的数据、从MP路由到数据路由总线的数据、从数据路由总线路由到PE的数据、以及从数据路由总线路由到MP的数据。例如，交换盒122.1可以被配置成为要从垫片存储器116、MP 120.1或两者传递到处理元件118.1的数据提供数据交换。此外，交换盒122.1可以被配置成将数据从垫片存储器116路由到MP 120.1。作为另一示例，交换盒122.2可以被配置成为要从处理元件118.1、MP 120.2和/或SB 122.1传递到处理元件118.2的数据提供数据交换。此外，交换盒122.2可以被配置成将数据从处理元件118.1路由到MP 120.2或SB 122.3，从SB 122.1路由到MP120.2或SB 122.3。在又一示例中，交换盒122.N可以被配置成为要从PE 118.N-1、MP120.N、SB 122.N-1或这三个源的任意组合传递到处理元件118.N的数据提供数据交换。此外，交换盒122.N可以被配置成在PE 118.N-1、MP 120.N、SB 122.N-1和垫片存储器116之间路由数据。SB也可以被称为数据交换单元(data switching unit)。

在一些实施例中，向量处理单元的输出端口(例如，每个MP、每个SB和每个PE)可以是向量地址或向量数据端口。输出端口处的地址或数据缓冲器可被视为向量寄存器。例如，耦合到SB 122.2的PE 118.1的一个输出端口处的数据缓冲器可以被视为向量寄存器，用于保持到SB122.2的输入值的向量。耦合到SB 122.2的PE 118.1的另一输出端口处的数据缓冲器可被视为另一向量寄存器，以保持到SB 122.2的输入值的另一向量。此外，耦合到SB122.2的SB 122.1的输出端口处的数据缓冲器可以被视为向量寄存器，用于保持要被传递到SB 122.2的数据值的向量。

在一个实施例中，向量处理单元的输出端口处的数据或地址缓冲器可以映射到被标记为VA、VB、IA、IB、IC、ID、IE、IF、IG、IH和DVA的向量寄存器。VA和VB可以是用于PE的输出数据缓冲器的向量寄存器。IA、IB、IC、ID、IE、IF、IG和IH可以是在SB的输出端口的输出数据缓冲器的向量寄存器，输出端口耦合到后继SB或垫圈存储器的输入端口。DVA可以是用于MP的输出地址缓冲器的向量寄存器。此外，虚拟向量寄存器MA和MB可以映射到从MP到SB的数据连接，使得SB可以将从存储器单元112读取的数据路由到PE的输入端口。MA和MB可以分别表示由共享存储器访问和由私有存储器访问获得的数据。DVA的宽度可以是KxG位。其它向量寄存器的宽度可以是KxM位。PE中的ALU可以被配置为M位(一个字)或2xM位(两个字)。为了支持2xM位操作，两个向量寄存器可以级联成寄存器对，并标记为VAB、IAB、ICD、IEF、IGH和MAB。例如，IAB可以指示级联的向量寄存器对(IB，IA)，其中，IB是较高的M位，IA是较低的M位。此处(，)表示来自两个向量寄存器的M位数据级联成的组件。

示例性数据路径可以通过SB 122.1到122.N的示例性内部连接来示出。例如，如图1所示，SB 122.1可以示出PE 118.1的两个输入可以耦合到来自MP 120.1的两个输出，SB122.2可以示出PE 118.2的两个输入可以耦合到来自MP 120.2的两个输出，并且PE118.2的另两个输入可以耦合到来自PE 118.1的两个输出，SB 122.3可以示出PE 118.3的两个输入可以耦合到来自MP 120.3的两个输出，并且PE 118.3的另两个输入可以耦合到来自PE 118.2的两个输出，依此类推，直到SB 122.N可以示出PE 118.N的两个输入可以耦合到来自MP120.N的两个输出,并且PE 118.N的另两个输入可以耦合到来自PE 118.N-1的两个输出。

为了简化措辞，MP(或MP 120)可以指MP 120.1-120.N中的一个，SB(或SB 122)可以指SB 122.1-122.N中的一个，并且PE(或PE 118)可以指PE 118.1-118.N中的一个。

可以为处理器100定义混合标量向量指令集(mixed-scalar-vector instructionset)。MP 120.1-120.N、SB 122.1-122.N和PE 118.1-118.N可以是处理器100的向量处理单元，并且序列发生器106可以是处理器100的标量处理单元。指令集可以被设计为使得一个指令可以由一个可编程单元执行。例如，在一个实施例中，每个指令可以具有32位，并且每个指令的某些位(例如，最高有效的4位、最低有效的4位或其他位)可以标识执行该指令的单元(例如，序列发生器106或PE阵列114的一个组件)。

内核程序可以由一系列指令组组成。标量指令可以生成向量指令中使用的参数，并管理环路和分支。向量指令可以配置列中的数据路径，控制数据流并执行数据处理操作。一个组的指令可以配置一个列。用于PE的包括无操作(no-operation，NOP)的指令可以是该组的分隔符。在一些实施例中，一个组中的指令可以被组织成使得为向量指令生成参数的标量指令被放置在向量指令之前。列不直接解码向量指令。相反，序列发生器106可以解码标量和向量指令，执行解码后的标量指令，并将解码后的向量指令打包到配置中，并将存储器端口(MP 120.1-120.N)、交换盒(SB 122.1-122.N)和处理元件(PE 118.1-118.N)的配置发送到列。

在各种实施例中，处理器100可以用于执行单指令多线程(single instructionmultiple threads，SIMT)执行。一组线程可以形成块(例如，线程块)，一组块可以被组织成网格(grid)。可以在内核程序执行之前为内核程序定义上述网格和块的组织。每个块和线程可以分别具有唯一的块和线程标识符(例如，块ID(block ID)和线程ID(thread ID))。三维网格中块的块ID可以计算为blockId＝blockIdx.x+(blockIdx.y*gridDim.x)+(blockIdx.z*(gridDim.x*gridDim.y))。变量blockIdx.x、blockIdx.y和blockIdx.z可以分别是块的x轴、y轴和z轴上的块ID。变量gridDim.x和gridDim.y可以分别是x轴和y轴上的网格大小。“*”运算符是乘法运算符。三维块中线程的线程ID可以计算为threadId＝blockId*(blockDim.x*blockDim.y*blockDim.z)+threadIdx.x+(threadIdx.y*blockDim.x)+(threadIdx.z*(blockDim.x*blockDim.y))。变量threadIdx.x、threadIdx.y和threadIdx.z可以分别是线程的x轴、y轴和z轴上的线程ID。变量blockDim.x、blockDim.y和blockDim.z可以分别是x轴、y轴和z轴上的块大小(block dimension)。

如本文所使用的，大写字母X、Y和Z可分别表示线程块在块的x轴、y轴和z轴上的大小。在一个实施例中，X、Y和XYZ的值(例如，X乘以Y乘以Z的乘积)可以由系统在启动内核之前在序列发生器106的本地存储器812中设置。并且序列发生器106可以从本地存储器812加载X、Y和XYZ，将X、Y和XYZ存储在序列发生器106的标量寄存器中，并且将X、Y和XYZ作为配置中的立即数传递给列(例如，列中的MP)。

图2示意性地示出了根据本公开的实施例的处理元件(PE)200。PE 200可以是PE118的实施例。PE 200可以包括算术逻辑单元(arithmetic logic unit，ALU)202、多个数据缓冲器(例如，D-FIFO 204.1和204.2)、计数器206、多个数据输出(例如，208.1和208.2)、多个数据输入(例如，210.1到210.6)、配置输入212和配置缓冲器(例如，C-FIFO 214)。在一个实施例中，ALU 202可以是一个ALU(例如，用于一次处理一个数据的一个ALU，并且可以被称为标量ALU)。在大多数实施例中，ALU 202可以是多个ALU(或称为向量ALU)，例如K个ALU，并且单指令多线程(single instruction multiple threads，SIMT)操作可以由PE执行。如本文所使用的，大写字母K可以被称为ALU的向量大小，并且示例K可以是32。应当注意，例如，对于向量寄存器和向量数据总线，相同的向量大小K可以应用于MP、SB中。请注意，标量ALU可以是向量ALU的特例，其中，向量大小为1。

从数据输入210.1到210.6接收的数据可以表示为A、B、C、D、E和F。发送到数据输出208.1和208.2的数据可以表示为VA和VB。在ALU 202可以是一个ALU的实施例中，数据输入210.1至210.6以及数据输出208.1和208.2的宽度可以是M位。通过配置，ALU的宽度可以被配置为M位(一个字)或2xM位(两个字)。如果宽度为M位，ALU的输入为A、B和C。ALU的输出为VA。如果宽度为2xM位，ALU的输入为(B，A)、(D，C)和(F，E)。ALU的输出为(VB，VA)。此处(，)表示M位数据的级联。例如，当M为8时，ALU的输入和输出可以是8位或16位；当M为16时，ALU的输入和输出可以是16位或32位；当M为32时，ALU的输入和输出可以是32位或64位；诸如此类。输入数据A、B、C、D、E和F以及输出数据VA和VB可以是M位。在ALU 202可以作为向量ALU的实施例中，数据输入210.1至210.6以及数据输出208.1和208.2可以是KxM位的向量。输入数据A、B、C、D、E和F以及输出数据VA和VB可以是KxM位的向量。

数据缓冲器204.1和204.2可以耦合到数据输出208.1和208.2以临时存储输出数据。可以分别映射到向量寄存器VA和VB的数据缓冲器204.1和204.2可以被用于将PE的时序与后续SB或垫片存储器的时序去耦。在一个实施例中，缓冲器可以被实现为FIFO(例如，用于数据缓冲器的D-FIFO，用于配置缓冲器的C-FIFO)。

配置缓冲器C-FIFO 214可以从配置输入212接收配置，并且在数据路径的任何执行开始之前存储接收到的配置，配置输入212可以经由配置总线外部耦合到序列发生器106。PE 200的配置可以称为PE配置。PE 200可以在处理线程块时被静态配置，例如，PE 200可以被配置中指定的指令编程以执行流水线的一个级。当线程块中的数据通过PE 200时，不能改变任何指令。配置参数XYZ中的一个可用于获得可以由ceil(XYZ/K)指定的执行次数。此处函数ceil(x)返回大于或等于x的最小整数值。计数器206可以被编程为执行次数，并用于对通过数据输出208.1的数据进行计数。当计数器值达到执行次数时，可以应用新的配置。因此，可以在每个PE中提供重构能力。在一个实施例中，指令的指定执行次数可以被称为NUM_EXEC，并发执行的K个线程可以称为warp，并且对于一个数据路径中的所有组件，该NUM_EXEC可以相等。例如，对于线程数TH＝1024个线程的线程块，NUM_EXEC＝ceil(1024/32)＝32。计数器206可以被称为warp计数器。

图3示意性地示出了根据本公开的实施例的处理器100的存储器单元300。存储器单元300可以是存储器单元112的实施例，并且可以用作PE阵列114的片上高速缓存。存储器单元300可包括多个存储器组(例如，表示为302.1的存储器组0、表示为302.2的存储器组1、表示为302.J的存储器组J-1等)、用于共享存储器访问的多个存储器高速缓存303.1至303.J以及用于私有存储器访问的多个存储器高速缓存304.1至304.J。高速缓存303.1至303.J以及高速缓存304.1至304.J也可以称为被高速缓存块。存储器组302.1至302.J中的每一个可以耦合到用于共享存储器访问的相应高速缓存303和用于私有存储器访问的相应高速缓存304。例如，存储器组302.1可以耦合到高速缓存303.1和高速缓存304.1，存储器组302.2可以耦合到高速缓存303.2和高速缓存304.2，等等，直到存储器组302.J可以耦合到高速缓存303.J和304.J。在一个实施例中，存储器单元300可以是J路交织存储器，其中，J可以是2的幂。

存储器单元300中的存储位置可以通过存储器地址来标识。在一个实施例中，存储器地址的某一段可以用于选择存储器组，并且该存储器地址的其余部分(可以被称为组地址(bank address))，可以用于标识存储器组中的位置。用于选择存储组的该段存储地址可以被称为组选择位(bank selection bit)并且具有L的位宽，在J为2的幂的实施例中，L的位宽可以等于J的位宽减1(例如，L是log₂(J))。因此，存储器地址可以指完整地址，而组地址可以指存储器地址中没有组选择位的部分。

在一些实施例中，存储器单元300可以是字可寻址的。也就是说，存储器单元300中的每个字可以由存储器地址唯一指向或标识，并且存储位置可以是可寻址的字位置。在一个示例(被称为存储器示例1)中，数字J可以是八(8)并且L可以是三，存储器地址的三(3)个最低有效位(least significant bit，LSB)可以用作组选择位，并且其他位可以用作组地址。如果存储组的位宽是64位，并且组地址指向的数据字是16位(假设PE 118中的ALU的字大小是16位)，则存储组中的64位可以保持4个16位数据字(word of data)。需要注意的是，每个存储组的组地址的2位可以是位置位，并且指示16位数据的字的位置(例如，组地址的位[1：0]指向位置3、2、1、或0的16位字)，并且组地址的其余部分可以是可以指向存储组中的64位数据的存储位置的地址。例如，十进制存储器地址15的二进制表示的最后四位是“1111”，并且其他位是零，三个LSB是“111”，两个位置位为“01”，并且该存储器地址的其余部分是零，使得该十进制存储器地址可以指向存储器组7中的地址0处的64位数据的第二个字。因此，地址可以指没有位置位的组地址部分。

在一些实施例中，存储器单元300可以是字节可寻址的。也就是说，存储器单元300中的每个字节可以由存储器地址唯一指向或标识，并且存储位置可以是可寻址的字节位置。作为示例，存储器地址的组选择位可用于选择存储器存储组，其他位可用作组地址。如果存储器组的位宽是(K/J)xM位，则存储器组中的数据可以保持(K/J)x(M/8)个字节。组地址的一段可用于选择字节位置，并且该组地址的其余部分可用于标识地址，该地址可以指向存储组中(K/J)xM位数据的存储位置。如果字大小是W(＝M/8)字节并且W大于1，则存储器地址可以包括数据大小位(data size bit)，该数据大小位的位宽等于W位宽-1(即log₂(W))。数据大小位可以以字节为单位指示存储器中的数据大小，并且当数据大小为W字节时，数据大小位可以全部为“0”。数据大小位的数据大小和位宽可以分别由DS和DSW表示。数据大小可以是2的幂。存储器地址必须能被DS整除。如果数据大小是W个字节(DS＝W)，则存储器地址的DSW位LSB为零，并且存储器地址的其余部分相当于字可寻址的存储器单元的存储器地址。当W大于DS时，对于数据加载过程，从存储器单元返回的数据的高位W-DS字节可以用零填充，而对于数据存储过程，发送到存储器单元的数据的高位W-DS字节可以被忽略。

在一个示例(被称为存储器示例二)中，对于字大小为2字节(W＝2且M＝16)、数据大小位是1位、存储器的位宽是64位(4个字)、并且存储器组的数量J为8的25位存储器地址，存储器地址的位[0]可以标识字中的字节(例如，高位字节或低位字节)，存储器地址的位[3：1]可以是组选择位并且标识8个存储组中的一个存储组，存储器地址的位[5：4]可以是位置位并且标识64位数据的4个字中的一个字，存储器地址的其余部分可以是指向存储器组中的64位数据的存储位置的地址。例如，如果DS为1，则存储器地址122可以指向组5(例如，位[3：1]＝5)中地址1(例如，位[24：6]＝1)处的字节6(例如，位[0]＝0，位[5：4]＝3)，或如果DS为2，则指向组5中的地址1处的字3(字节7和字节6)。

需要注意的是，J是8和存储器组宽度是64只是示例，并且在各种实施例中，J和存储器组宽度可以是其他数字。另外，用于组选择的位数可以取决于存储器单元300中的存储组的数量。当J是8时，可以使用3位进行组选择，因为3位的取值范围是零(0)至七(7)。

每个高速缓存303可以单独地耦合到所有多个MI 124.1-124.N，用于经由连接端口306进行共享存储器访问，并且每个高速缓存304可以单独地耦合到所有多个MI124.1-124.N，用于经由连接端口308进行私有存储器访问。连接端口306和308中的每一个可以使用两个订阅(subscription)来标识其连接，其中，第一订阅标识存储器高速缓存(通过存储器高速缓存订阅1到J)，以及第二订阅标识MI(通过MI订阅1到N)。例如，连接端口306.1.1可以用于存储器高速缓存303.1和MI 124.1的共享存储器访问，连接端口306.2.1可以用于存储器高速缓存303.2和MI 124.1的共享存储器访问，依此类推，直到连接端口306.J.1可以用于存储器高速缓存303.J和MI 124.1的共享存储器访问；连接端口306.1.2可以用于存储器高速缓存303.1和MI 124.2的共享存储器访问，连接端口306.2.2可以用于存储器高速缓存303.2和MI 124.2的共享存储器访问，依此类推，直到连接端口306.J.2可以用于存储器高速缓存303.J和MI 124.2的共享存储器访问；连接端口306.1.N可以用于存储器高速缓存303.1和MI 124.N的共享存储器访问，连接端口306.2.N可以用于存储器高速缓存303.2和MI 124.N的共享存储器访问，依此类推，直到连接端口306.J.N可以用于存储器高速缓存303.J和MI 124.N的共享存储器访问。

类似地，连接端口308.1.1可以用于存储器高速缓存304.1和MI 124.1的私有存储器访问，连接端口308.2.1可以用于存储器高速缓存304.2和MI 124.1的私有存储器访问，依此类推，直到连接端口308.J.1可以用于存储器高速缓存304.J和MI 124.1的私有存储器访问；连接端口308.1.2可以用于存储器高速缓存304.1和MI 124.2的私有存储器访问，连接端口308.2.2可以用于存储器高速缓存304.2和MI 124.2的私有存储器访问，依此类推，直到连接端口308.J.2可以用于存储器高速缓存304.J和MI 124.2的私有存储器访问；连接端口308.1.N可以用于存储器高速缓存304.1和MI 124.N的私有存储器访问，连接端口308.2.N可以用于存储器高速缓存304.2和MI 124.N的私有存储器访问，依此类推，直到连接端口308.J.N可以用于存储器高速缓存304.J和MI 124.N的私有存储器访问。

应当注意，高速缓存303的数量和高速缓存304的数量两者都可以与由大写字母J表示的存储器组的数量相匹配。并且MI 124的数量可以与由大写字母N表示的列的数量相匹配。存储器组的数量不需要与向量大小相同。例如，向量(例如，向量ALU、向量地址、向量数据端口)可以具有向量大小K，PE阵列可以具有列的数量N，并且存储器单元可以具有存储器组的数量J。并且K、N和J可以都不同。在一个实施例中，K可以被J整除，J可以是2的幂，并且J的位宽减1可以是L(例如，L是log₂(J))。例如，J和L可以分别是八(8)和三(3)，K可以是32，N也可以是32。

共享存储器模式可以允许随机存储器访问。假设数据路径的向量大小是K。那么共享存储器地址总线的位宽可以是KxG，其中G是地址空间的位宽(例如，存储器示例2中的25)。存储器地址可以通过使用向量寄存器来指定。用于在共享存储器模式下加载数据的汇编代码示例是加载指令，例如，要由MP执行的LOAD[DVA+VA]、DEST，其中DVA、VA、以及DEST可以是向量寄存器。DVA可以包含K个基址，并且VA可以提供K个地址偏移(address offset)。DVA和VA可以分别由前一列中的MP和PE更新。DVA和VA的相加可以由MP中的向量ALU(例如，MP 120.2的ALU 718)来执行。DEST可以是从存储器单元300加载的数据的K个字的目的地，DEST可以是同一列中SB的输出数据端口(例如，SB 122.2的数据端口526.1至526.8中的任意一个)或同一列中SB的输入数据端口MA(例如，SB 120.2的502.1)，数据的K个字可以被转发到同一列中的PE(例如，经由SB 120.2的506.1的PE 118.2的210.1)。如果K个存储器地址中的一些地址被定向到同一个存储器组，存储器接口可以解决争用，并分配多个周期以从存储器组加载。

例如，如果加载(LOAD)指令将由MP 120.2执行，则DVA和VA可以分别由前一列中的MP 120.1和PE 118.1更新。DVA可以是映射到前一列中的MP的输出地址缓冲器(例如，MP120.1的输出地址缓冲器720)的向量寄存器，并且VA可以是映射到前一列中的PE的输出数据缓冲器(例如，PE 118.1的输出数据缓冲器204.1)的向量寄存器。DEST可以是映射到与执行加载(LOAD)指令的MP在同一列中的SB的输出数据缓冲器的向量寄存器IA至IH中的一个(例如，SB 122.2的输出数据缓冲器522.7至522.14中的一个)，或是从MP到与MP在同一列中的SB映射到共享存储器端口的读数据连接的虚拟向量寄存器MA(例如，MP 120.2的端口728和SB 122.2的端口502.1之间的连接)。

在一些实施例中，在私有存储器访问模式中，并行线程的数据可以以连续存储器地址存储在存储器单元300中。需要注意的是，尽管存储器地址可以是连续的，但是由存储器地址指向的存储位置实际上可以分布在多个存储器组中。例如，如果PE 118中有32个并行执行32个线程(例如warp)的ALU，并且每个ALU用于对16位数据字进行操作，则向量数据总线的大小可以是32乘以16位(共512位)。对于存储器示例1，因为每个存储器地址的三个LSB选择一个存储器组，所以连续存储器地址可以指向不同的存储器组。例如，如果第一个线程(例如线程0)的数据字存储在存储器地址15，则32个数据字的连续存储器地址的范围可以是15到46。存储器地址15的字可以位于组7的地址0处的64位数据的位置1(例如，位[31：16])。下一个线程(例如线程1)的字可以存储在存储器地址16，该存储器地址16的2位置位是“10”且3位LSB是“000”，并且可以指向位于组0的地址0的64位数据的位置2(例如，位[47：32])处的字。并且下一个线程(例如线程2)的字可以存储在存储器地址17，该存储器地址17的2位置位是“10”且3位LSB是“001”，可以指向位于组1的地址0的64位数据的位置2(例如，位[47：32])处的字。可以推断的是，在组7的地址0处的64位数据的位置2和3处的其他字可以用于存储器地址23和31，上述存储器地址23和31可以分别是第9和第17线程(例如，线程8和16)的字，并且在组7的地址1处的64位数据的位置0处的字可以用于存储器地址39和第25线程(例如，线程24)。并且在存储组0的地址0处的64位数据的位置3处的另一个字可以用于存储器地址24和第10线程(例如，线程9)。另外，在存储组0的地址1处的64位数据的位置0和1处的字可以分别用于存储器地址32和40，以及分别用于第18和第26线程(例如，线程17和25)。

指向存储器组的存储器地址可以与指向同一存储器组的另外3个存储器地址相关联，并且可以访问64位数据(4个数据字)，该64位数据可以由4个存储器地址标识，并存储在1个存储器位置或2个连续存储器位置中。如果四个字中的第一个字位于存储位置的开头(例如，在存储位置的第一个字处)，则上述4个数据字可以位于一个存储位置中，或如果四个字中的第一个字没有位于存储位置的开头(例如，在存储位置的第二个字、第三个字、或第四个字处)，则上述4个数据字可以位于两个连续的存储位置中。例如，存储器地址15、23、31和39可以被定向到相同的存储器组7。因此，存储器地址15可以与存储器地址23、31和39相关联，并用于访问4个数据字，一个字用于线程0，又一个字用于线程8，另一个字用于线程16，还有一个字用于线程24。被定向到组7的存储器地址15可以分别指向存储器地址15的地址0处的存储器位置的位置1(例如，位[31：16])处的1个字、存储器地址23的地址0处的存储器位置的位置2(例如，位[47：32])处的1个字、存储器地址31的地址0处的存储器位置的位置3(例如，位[63：48])处的1个字、以及存储器地址39的地址1处的存储器位置的位置0(例如，位[15：0])处的1个字。存储器接口可以从存储器地址15中移除组选择位，并向存储器组7提供组地址1。

被定向到存储器组0到7的存储器地址可以分别是16、17、18、19、20、21、22、以及15。私有存储器访问接口可以从一个标量存储器地址生成8个组地址。提供给存储组0到7的组地址分别是2、2、2、2、2、2、2、以及1。存储器地址15的64位数据位[47：0]可以存储在组7的地址0处的数据位[63：16]中，并且位[63：48]可以存储在组7的地址1处的数据位[15：0]中。存储器地址16的64位数据位[31：0]可以存储在组0的地址0处的数据位[63：32]中，并且位[63：32]可以存储在组0的地址1处的数据位[31：0]中。类似地，存储器地址17至22的64位数据的位[31：0]可以存储在存储组1至6的地址0处的数据位[63：32]中，并且位[63：32]可以存储在存储组1至6的地址1处的数据位[31：0]中。

因为来自连续存储器地址的16位数据字可以在8个存储器组之间交织，所以对于数据加载过程操作，存储器接口124可以从每个存储器组收集64位数据，将每个64位数据拆分为4个16位数据字，并对8组这样的4个16位数据字进行解交织，以通过将来自存储器组7、组0、组1、…、以及组6的数据位[15：0]分别放置到512位向量的位[15：0]、位[31：16]、位[47：32]、…、以及位[127：112]，然后将来自存储器组7、组0、组1、…、以及组6的数据位[32：16]放置到位[143：128]、位[159：144]、位[175：160]、以及位[255：240]，等等，从而组织用于输出的512位向量。

对于存储器示例2，字大小可以是2字节(W＝2)，并且向量数据总线的大小可以是32乘以16位(共512位)。为从存储器单元加载/向存储器单元存储一个512位的向量数据，如果DS是1，则可以有32个连续存储器地址，或如果DS是2，则可以有32个连续偶数存储器地址。

在DS为2的示例中，如果第一线程的数据的两个字节位于存储器地址122，则所有512位数据的对应存储器地址的范围是从122到184(例如，甚至是连续的存储器地址)。指向存储器组0到7的存储器地址分别是128、130、132、134、136、122、124、以及126。提供给存储组0到7的组地址分别是16、16、16、16、16、14、14、以及14。存储512位数据的地址和组是：对于位[399：384]、位[271：256]、以及位[143：128]，组5的地址2的数据位[47：0]；对于位[15：0]，组5的地址1的数据位[63：48]；对于位[415：400]、位[287：272]、以及位[159：144]，组6的地址2的数据位[47：0]；对于位[31：16]，组6的地址1的数据位[63：48]；对于位[431：416]、位[303：288]、以及位[175：160]，组7的地址2的数据位[47：0]；对于位[47：32]，组7的地址1的数据位[63：48]；对于位[447：432]、位[319：304]、位[191：176]、以及位[63：48]，组0的地址2的4个数据字；对于位[463：448]、位[335：320]、位[207：192]、以及位[79：64]，组1的地址2的4个数据字；对于位[479：464]、位[351：336]、位[223：208]、以及位[95：80]，组2的地址2的4个字的数据；对于位[495：480]、位[367：352]、位[239：224]、以及位[111：96]，组3的地址2的4个数据字；对于位[511：496]、位[383：368]、位[255：240]、以及位[127：112]，组4的地址2的4个数据字。

在DS为1的另一示例中，如果第一线程的数据的一个字节位于存储器地址122，则所有512位数据的对应存储器地址的范围是从122到153的连续存储器地址。指向存储器组0到7的存储器地址分别是128、130、132、134、136、122、124、以及126。提供给存储组0到7的组地址分别是16、16、16、16、16、14、14、以及14。存储512位数据的地址和组是：对于位[391：384]和位[263：256]，组5的地址2的数据位[15：0]；对于位[135：128]和位[7：0]，组5的地址1的数据位[63：48]；对于位[407：400]和位[279：272]，组6的地址2的数据位[15：0]；对于位[151：144]和位[23:16]，组6的地址1的数据位[63：48]；对于位[423：416]和位[295：288]，组7的地址2的数据位[15：0]；对于位[167：160]和位[39：32]，组7的地址1的数据位[63：48]；对于位[439：432]、位[311：304]、位[183：176]、以及位[55：48]，组0的地址2的数据位[31：0]；对于位[455：448]、位[327：320]、位[199：192]、以及位[71：64]，组1的地址2的数据位[31：0]；对于位[471：464]、位[343：336]、位[215：208]和位[87：80]，组2的地址2的数据位[31：0]；对于位[487：480]、位[359：352]、位[231：224]和位[103：96]，组3的地址2的数据位[31：0]；对于位[503：496]、位[375：368]、位[247：240]和位[119：112]，组4的地址2的数据位[31：0]。对于数据加载过程，32个数据字的所有高位字节(位[511,504]、位[495,488]、……、位[31，24]、以及位[15，8])可以被设置为0，或者对于数据存储过程，上述32个数据字的所有高位字节可以被忽略。

图4A示意性地示出了根据本公开的实施例的存储器接口(memory interface，MI)400。MI 400可以是图1的MI 124的实施例。在一些实施例中，存储器接口可以被称为根盒(root box)。MI 400可以包括用于共享存储器访问模式的共享存储器访问接口402和用于私有存储器访问模式的私有存储器访问接口403。私有存储器访问接口403可以包括耦合到存储器端口的地址端口404、写数据(write data，WData)端口406和读数据(read data，RData)端口408。私有存储器访问接口403还可以包括耦合到存储器单元300的多个地址端口410.1-410.J、多个WData端口412.1-412.J和多个RData端口414.1-414.J。共享存储器访问接口402可包括耦合到存储器端口的地址端口428、WData端口430和RData端口432。共享存储器访问接口402还可以包括耦合到存储器单元300的多个地址端口426.1-426.J、多个WData端口422.1-422.J和多个RData端口424.1-424.J。

对于到存储器单元300的连接，一组地址、WData和RData总线可以耦合到图3所示的一个连接端口306.1.1-306.J.N和308.1.1-308.J.N。例如，MI 124.1的地址端口410.1、WData端口412.1和RData端口414.1可以耦合到存储器单元300的连接端口308.1.1；MI124.1的地址端口410.J、WData端口412.J和RData端口414.J可以耦合到连接端口308.J.1。同时，MI 124.N的地址端口410.1、WData端口412.1和RData端口414.1可以耦合到存储器单元300的连接端口308.1.N；MI124.N地址端口410.J、WData端口412.J和RData端口414.J可以耦合到连接端口308.J.N。

类似地，MI 124.1的地址端口426.1、WData端口422.1和RData端口424.1可以耦合到存储器单元300的连接端口306.1.1；MI 124.1的地址端口426.J、WData端口422.J和RData端口424.J可以耦合到连接端口306.J.1。同时，MI 124.N的地址端口426.1、WData端口422.1和RData端口424.1可以耦合到存储器单元300的连接端口306.1.N；MI 124.N的地址端口426.J、WData端口422.J和RData端口424.J可以耦合到连接端口306.J.N。

在一个实施例中，耦合到存储器端口的WData端口和RData端口中的每一个可以被配置用于向量数据连接。例如，WData端口406可以是KxM位输入端口，RData端口408可以是KxM位输出端口。

在一些实施例中，地址端口404可以被配置为向量地址。但是并非所有实施例都需要将地址端口404配置为向量地址。在私有存储器访问模式中，一个warp的一个向量地址中的K个存储器地址可以根据线程ID以升序连续排列。因此，在一个实施例中，可能只需要将warp的所有线程中具有最小线程ID的线程的标量存储器地址从MP发送到MI。因此，在这样的实施例中，地址端口404可以用于标量存储器地址，并且地址端口404的宽度可以是G位。

另外，假设J小于或等于K，则每个组(例如，412和414)的数据端口的宽度可以是(K/J)xM位。由于存储器单元300可以是J路交织存储器，存储器地址的L位组选择可以确定在其中存储器地址的数据可以驻留的存储器组。这里L可以是J的位宽减去1。存储器示例1和存储器示例2的组选择位可以分别是存储器地址的位[2：0]和位[3：1]。一个向量数据中的所有K个数据可以均匀地分布在存储器组中，并且是可以在没有存储器争用的情况下进行访问。

在共享存储器访问模式中，一个向量地址中的K个存储地址可以彼此不同。由向量地址访问的数据可以被随机地分布在所有存储器组中，这可能导致存储器争用。地址端口428的宽度可以是KxG位。每个组(例如，422和424)的数据端口的宽度可以是M位。共享存储器访问接口402可以解决存储器争用。

图4B示意性地示出了根据本公开实施例的用于私有存储器访问的私有存储器访问接口403。在WData端口406和RData端口408的向量数据的宽度可以是KxM位。向量数据可以被分成K个M位数据的字，这K个字可以根据K个字的相应线程ID的顺序从0到K-1被索引。在地址端口404接收的标量存储器地址的宽度可以是G位。耦合至地址端口404的存储器总线可以被称为私有存储器地址总线。在存储器单元是字节可寻址存储器单元的实施例中，私有存储器地址总线还可以包括数据大小位(未在图4B中示出)。也就是说，当存储器单元112是字节可寻址存储器单元实施例时，地址端口404可以是G+DSW位，其中DSW是数据大小位的位宽。如果存储器单元是字可寻址的，则可以通过在地址端口404将字的索引(例如，0到K-1)加到标量存储器地址来计算K个字中的每一个的存储器地址；或如果存储器单元是字节可寻址的，则通过在地址端口404将乘以DS的索引加到标量存储器地址来计算K个字中的每一个的存储器地址，其中DS是以字节为单位的数据大小。

私有存储器访问接口403的组地址生成单元440可以生成K个存储器地址，检查K个存储器地址中的每一个的L位组选择位，并根据由组选择位指定的存储组索引对上述组地址进行分组。在每个分组中，可以选择最小的存储器地址，然后可以删除该存储器地址的组选择位。剩余的G减去L(G-L)位和字节可寻址存储器单元的数据大小位可以被分配给地址A_0到地址A_J-1中的一个，地址A_0到地址A_J-1的索引与该分组的组索引相匹配。

对于数据存储过程，私有存储器访问接口403的写数据拆分单元442可以将在WData端口406接收的KxM比位数据分成K个M位数据的字，并根据K个字的相应线程ID的顺序从0到K-1对这K个字进行索引。组地址生成单元440可以计算K个字的存储器地址，并根据由组选择位指定的存储组索引对上述组地址进行分组。地址属于同一分组的字可以被发送到私有存储器访问接口403的写数据组合单元446.1至446.J中的一个，字的索引与该分组的存储组索引匹配。每个写数据组合单元446按照字的索引的顺序组合分组中的字。每个分组的组合数据可以被分配给写数据WD_0至WD_J-1中的一个，并经由输出端口412.1至412.J中的一个被传送到存储器组。对于数据加载过程，可以经由输入端口414.1至414.J从存储器组接收(K/J)xM位读数据RD_0至RD_J-1。每个读数据可以被分配给分组的存储组索引与读数据端口的索引匹配的分组，然后由私有存储器访问接口403的读数据拆分单元448.1至448.J中的一个分成K/J个M位数据的字。可以从组地址生成单元440中生成的字的地址索引中检索分组中每个字的索引。来自所有J个分组的K个字可以按照每个字的索引的顺序重新排列，然后K个字可以由私有存储器访问接口403的读数据组合单元444组合。KxM位数据向量可以形成并被分配给RData端口408。

图4C示意性地示出了根据本公开的实施例的用于共享存储器访问的共享存储器访问接口402。地址端口428可以被配置用于KxG位的向量地址，例如A_0、A_1…A_K-1的K个G位地址。数字K可以对应于PE 118中的向量ALU的向量大小K。K个地址可以被传递到多个地址选择单元(例如，“选择2”单元416.1到416.J)。每个地址选择单元416.1到416.J可以将每个存储器组的索引作为输入，例如，“组0”的索引0……和“组J-1”的索引J-1，扫描从A_0到A_K-1的地址，拾取其组选择位与组索引匹配的所有地址，并通过地址端口426(例如，用于存储器组0高速缓存303.1的地址端口426.1，用于存储器组J-1高速缓存303.J的地址端口426.J，等等)将地址的剩余G-L位逐个发送到存储器组。地址端口426的宽度可以是G-L位。

因为多于一个的地址可以被定向到相同的存储器组，所以可以提供写数据选择单元(例如，“选择2”单元418.1到418.J)和读数据选择单元(例如，“选择”单元420.1到420.J)，以将正在写入存储器组或从存储器组读取的数据与发送到存储器组的地址相匹配。写数据选择单元418.1至418.J中的每一个可以从对应的地址选择单元416.1至416.J接收发送到对应的地址端口426.1至426.J的每个地址的索引(例如，0至K-1)，并且将具有相同索引的写数据(例如，WD_0至WD_K-1)中的一个发送到WData端口422.1至422.J(例如，用于存储器组0高速缓存303.1的WData端口422.1，用于存储器组J-1高速缓存303.J的WData端口422.J，等等)。例如，如果地址选择单元416.1将A_2、A_15和A_28的G-L位发送到地址端口426.1，则写数据选择单元418.1接收索引2、15和28，并将WD_2、WD_15和WD_28发送到WData端口422.1。读数据选择单元420.1至420.J中的每一个可以从对应的地址选择单元416.1至416.J接收发送到对应的地址端口426.1至426.J的每个地址的索引(例如，0至K-1)，并且将从RData端口424.1至424.J(例如，用于存储器组0高速缓存303.1的RData端口424.1，用于存储器组J-1高速缓存303.J的RData端口424.J，等等)接收的数据分配给具有相同索引的读数据中的一个(例如，RD_0至RD_K-1)。例如，如果地址选择单元416.1向地址端口426.1发送A_2、A_15和A_28的G-L位，则读数据选择单元420.1接收索引2、15和28，并将从RData端口424.1接收的数据分配给RD_2、RD_15和RD_28。

图5示意性地示出了根据本公开的实施例的交换盒(switch box，SB)500。SB 500可以是SB 122的实施例，并且可以包括多个数据输入和数据输出，以及将数据输入耦合到数据输出以进行数据交换的互连。SB 500的数据输入可以包括数据输入502.1、502.2、514.1、514.2和524.1至524.8。SB 500的数据输出可以包括数据输出504.1至504.3、506.1、506.2、508.1、508.2、510.1、510.2和526.1至526.8。

在外部，数据输入502.1和502.2可以耦合到MP的数据输出(例如，读数据端口)，并分别映射到虚拟向量寄存器MA和MB。数据输入502.1和502.2中的一个可以耦合到私有存储器访问数据输出，而另一个可以耦合到共享存储器访问数据输出。数据输出504.1可以耦合到MP的数据输入端口。数据输出504.2和504.3可以耦合到MP的数据输入(例如，写数据端口)。数据输出504.2和504.3中的一个可以耦合到私有存储器访问写数据端口，而另一个可以耦合到共享存储器访问写数据端口。数据输入514.1和514.2可以分别耦合到PE的数据输出208.1和208.2(例如，标记为VA和VB)(或者在SB 122.1的情况下垫片存储器的对应输出)。数据输入524.1到524.8可以分别耦合到前续列的SB的数据输出526.1到526.8(或者在SB 122.1的情况下垫片存储器的对应输出)。506.1、506.2、508.1、508.2、510.1、510.2的数据输出可以分别耦合到PE的数据输入端口210.1到210.6。来自数据输出506.1、506.2、508.1、508.2、510.1、以及510.2的数据可以表示为A、B、C、D、E、以及F，并且来自数据输入514.1和514.2的数据可以表示为VA和VB。A、B、C、D、E、以及F可以是如本文所描述的到PE118的输入数据，并且VA和VB可以是来自PE 118的输出数据。

SB 500还可以包括配置缓冲器518和对应的配置输入516。配置缓冲器518可以被实现为先进先出缓冲器，并被称为C-FIFO 518。配置输入516可以外部耦合到配置总线，该配置总线耦合到序列发生器106，以便SB 500从序列发生器106接收配置。SB 500的配置可以称为SB配置。此外，SB 500还可以包括多个计数器520.1-520.27。除了数据输入502.1和502.2之外，其他数据输入和所有数据输出中的每一个都可以具有对应的计数器520。此外，SB 500还可以包括多个数据缓冲器522.1-522.14，其可以被实现为数据先进先出缓冲器并被称为D-FIFO 522.1-522.14。D-FIFO 522.1-522.14中的每一个可以为数据输出506.1、506.2、508.1、508.2、510.1、510.2、以及526.1-526.8中的每一个提供相应的输出缓冲器。D-FIFO522.7-522.14可以分别映射到向量寄存器IA、IB、IC、ID、IE、IF、IG、以及IH。

在SB 500内部，数据输入502.1可以耦合到数据输出506.1、506.2、508.1、510.1、以及526.1至526.8。数据输入502.2可以耦合到数据输出506.1、506.2、508.1、508.2、510.2、以及526.1至526.8。数据输入514.1可以耦合到数据输出504.1-504.3、506.1、506.2、508.1、510.1、以及526.1至526.8。数据输入514.2可以耦合到数据输出504.2、504.3、506.1、506.2、508.1、508.2、510.2、以及526.1至526.8。数据输入524.1、524.3、524.5、以及524.7中的每一个可以耦合到数据输出504.1-504.3、506.1、506.2、508.1、510.1以及输出526.1、526.3、526.5和526.7中的对应输出。数据输入524.2、524.4、524.6和524.8中的每一个可以耦合到数据输出504.1-504.3、506.1、506.2、508.1、508.2、510.2以及输出526.2、526.4、526.6和526.8中的对应输出。例如，数据输入524.1可以耦合到504.1-504.3、506.1、506.2、508.1、510.1和526.1，数据输入524.2可以耦合到504.1-504.3、506.1、506.2、508.1、508.2、510.2和526.2等。应该注意的是，在SB 500内部，输入和输出之间的耦合可以基于在SB 500处应用的当前配置而开启(例如，连接)或关闭(例如，断开)。此外，PE的D、E和F端口可以用于2xM位配置。只有寄存器对的较高M位(例如，VB、IB、ID、IF、IH和MB)可以被分配给D和F，并且只有寄存器对的较低M位(例如，VA、IA、IC、IE、IG和MA)可以被分配给E。

在数据端口处的每个计数器520.1-520.27可以独立地负责对通过数据端口的数据进行计数。当可以将一个或多个配置加载到C-FIFO 518中时，每个配置可以指定执行次数(例如，NUM_EXEC)。这些计数器可以被称为warp计数器。在一个配置的执行期间，所有计数器520.1-520.27可以独立地对通过数据端口的数据数量进行计数。当所有warp计数器达到配置中指定的执行次数时，可以应用C-FIFO 518中的下一配置。

使用warp计数器的类似方法可以应用于PE 118内部和存储器端口120。因为这些计数器可以有助于配置和重配置具有这样的计数器的每个组件，所以这些计数器可以被称为可重构计数器，并且具有这样的计数器的组件可以被称为可重构单元。处理器100的一个实施例可以使用各种可重配置单元提供大规模并行数据处理，并且可以被称为可重构并行处理器(reconfigurable parallel processor，RPP)。

图6示意性地示出了根据本公开的实施例的垫片存储器600。垫片存储器600可以是图1所示的垫片存储器116的实施例。垫片存储器600可以包括多个用于临时存储数据的缓冲器和一个用于地址的缓冲器。数据缓冲器可以被实现为先进先出(FIFO)缓冲器，并被称为D-FIFO(例如，D-FIFO 602.1-602.10)。地址缓冲器可以被实现为地址FIFO(例如，FIFO601)。此外，垫片存储器600可包括多个数据输入(例如614、616.1-616.2和618.1-618.8)和多个数据输出(例如608、610.1-610.2和612.1-612.8)。

输入614可以耦合到MP 120.N的输出，并且输出608可以耦合到MP 120.1的输入。在垫片存储器600内部，A-FIFO 601可以耦合在输入614和输出608之间。输入616.1和616.2可以分别耦合到PE 118.N的输出208.1和208.2。输出610.1和610.2可以耦合到SB122.1的输入514.1和514.2。在垫片存储器600内部，D-FIFO 602.1可以耦合在输入616.1和输出610.1之间，D-FIFO 602.2可以耦合在输入616.2和输出610.2之间。输入618.1-618.8可以分别耦合到SB 122.N的输出526.1-526.8。输出612.1-612.8可以分别耦合到SB 122.1的输入524.1-524.8。在垫片存储器600内部，D-FIFO 602.3-602.10可以分别耦合在输入618.1-618.8和输出612.1-612.8之间。

图7示意性地示出了根据本公开的实施例的处理器的存储器端口(MP)700。MP700可以是MP 120的实施例。存储器端口700可以包括地址端口702、WData端口704和RData端口706，其分别耦合到共享存储器访问接口402的地址端口428、WData端口430和RData端口432；地址端口708、WData端口710和RData端口712，分别耦合到私有存储器访问接口403的地址端口404、WData端口406和RData端口408；数据端口724、WData端口726、RData端口728、另一WData端口730和另一RData端口732，分别耦合到SB 500的数据端口504.1、504.2、504.3、502.1、以及502.2。

存储器端口700还可以包括配置输入734和配置缓冲器(例如，C-FIFO)736。MP配置可以包括要在MP上执行的指令，例如，从存储器单元加载数据和将数据存储到存储器单元的LOAD和STORE指令。存储器端口700还可以包括地址输入端口716、ALU 718、地址缓冲器(例如，A-FIFO)720和地址输出端口722。地址输入端口716可以耦合到前续列的MP的地址输出端口722(或者在MP 120.1的情况下垫片存储器600的地址输出端口608)，并且地址输出端口722可以耦合到后续列的MP的地址输入端口716(或者在MP 120.N的情况下垫片存储器600的地址输入端口614)。ALU 718可以对从地址端口716接收的地址和从数据端口724接收的数据执行操作，并将结果地址输出到地址端口702。此外，ALU 718可以将结果地址输出到地址端口722，或者将从地址端口716接收的地址传递到地址端口722。在地址从地址端口722输出之前，A-FIFO 720可以暂时存储来自ALU 718的地址。A-FIFO 720可以被映射到向量寄存器DVA。

存储器端口700还可以包括地址计算单元714。地址计算单元714可以用于使用线程块大小信息和寻址参数(例如基址、以及y轴和z轴上的存储器地址的间隙(gap))来生成用于私有存储器访问的存储器地址。基址可以是线程的块(或网格)的数据的起始存储器地址。线程块大小信息和寻址参数可以在MP配置中被传递到MP 700。

存储器端口700还可以包括多个用于计算一个配置的执行次数的计数器740.1-740.8(例如，NUM_EXEC)。每个计数器740可以与数据端口或地址端口相关联。计数器740.1-740.8中的每个计数器可以独立地负责对通过相关联端口的数据进行计数。例如，计数器740.1可以与地址端口702相关联，计数器740.2可以与WData端口704相关联，计数器740.3可以与地址端口708相关联，计数器740.4可以与WData端口710相关联，计数器740.5可以与地址输入端口716相关联，计数器740.6可以与地址输出端口722相关联，计数器740.7可以与RData端口728相关联，以及计数器740.8可以与RData端口732相关联。计数器740.1-740.8可以是warp计数器。在执行一个配置期间，所有计数器740.1-740.8都可以独立地计算通过数据端口的数据数量。当所有warp计数器达到配置中指定的执行次数时，可以应用C-FIFO 736中的下一配置。

在一些实施例中，私有存储器模式可以允许使用标量存储器地址进行顺序存储器访问。因为K个数据字(例如，warp的K个并发线程)的存储器地址可以是连续的，所以私有存储器访问模式中的地址总线可以只需要提供K个数据字的第一个元素的存储器地址。例如，在存储器示例1和存储器示例2中，只有第一数据的存储器地址可以由MP提供(例如，存储器示例1中的15和存储器示例2中的122)，存储器接口400和存储器单元300可以确定所有32个线程的其余存储器地址，并加载(或存储)所有32个线程的数据(例如，512位数据)。在这些实施例中，因为从MP向存储器接口400仅提供标量存储器地址，所以地址总线的位宽可以只需要是G，而不是像在共享存储器访问模式中位宽为KxG。如果存储器单元112是字节可寻址存储器单元，则MP可以在配置中接收数据大小DS作为立即值，并经由标量存储器地址总线的数据大小位将立即值发送到相应的存储器接口。

在私有存储器模式中用于从存储器单元112加载数据的汇编代码的示例可以是：LOAD％BASE,％STRIDE_Y,％STRIDE_Z,DEST(如果存储器单元112是字可寻址存储器单元(word addressable memory unit))；以及LOAD DS,％BASE,％STRIDE_Y,％STRIDE_Z,DEST(如果存储器单元112是字节可寻址存储器单元(byte addressable memory unit))。该加载(LOAD)指令可以由MP执行，以将warp的数据的K个字加载到DEST向量寄存器。该执行可以重复NUM_EXEC次以覆盖整个线程块(例如，warp的NUM_EXEC)。DEST向量寄存器可以是映射到从MP的输出端口732到同一列的SB的输入端口502.2的读数据连接的虚拟向量寄存器MB，或是映射到与MP同一列中的SB的输出数据缓冲器的向量寄存器IA至IH中的一个。寄存器MB可以向与MP在同一列的PE发送数据。寄存器IA到IH可以向后一列的SB发送数据。DS可以是以字节为单位表示数据大小的立即值。

如本文所用，参数名称前面的％符号可以指被指定以存储该参数值的标量寄存器。加载(LOAD)指令中的BASE、STRIDE_Y、以及STRIDE_Z可以由序列发生器106中的标量寄存器提供。BASE、STRIDE_Y、STRIDE_Z、以及DS可以称为寻址参数，并且作为从序列发生器106传送到MP的立即值嵌入到配置中。BASE可以是要处理的数据的起始存储器地址。当存储器单元112是字可寻址存储器单元时，STRIDE_Y和STRIDE_Z可以分别是y轴和z轴上的存储器地址的间隙，或当存储器单元112是字节可寻址存储器单元时，STRIDE_Y和STRIDE_Z可以分别是y轴和z轴上的存储器地址的间隙除以DS。STRIDE_Y和STRIDE_Z可以分别通过blockDim.x*gridDim.x和(blockDim.x*gridDim.x)*(blockDim.y*gridDim.y)给出，其中“*”是乘法运算符。

在私有存储器模式中用于将数据存储到存储器单元112的汇编代码的示例可以是：SRC,％BASE,％STRIDE_Y,％STRIDE_Z(当存储器单元112是字可寻址存储器单元时)，以及STORE DS,SRC,％BASE,％STRIDE_Y,％STRIDE_Z(当存储器单元112是字节可寻址存储器单元时)。该存储(STORE)指令可以由MP执行，以存储来自向量寄存器SRC的warp的数据的K个字。该执行可以重复NUM_EXEC次以覆盖整个线程块(例如，warp的NUM_EXEC)。SRC向量寄存器可以是映射到PE和前一列的SB的输出数据缓冲器的向量寄存器VA、VB、以及IA到IH中的一个。SRC向量寄存器可以将数据发送到同一列的SB的输出端口504.3，该输出端口可以耦合至同一列的MP的输入端口730。DS可以是以字节为单位表示数据大小的立即值。寻址参数BASE、STRIDE_Y、STRIDE_Z、以及DS可以以与加载(LOAD)指令相同的方式作为立即值嵌入到配置中。通过使用相同寻址参数的存储(STORE)和加载(LOAD)指令来生成的地址可以是相同的。

如果数据是720P(1280×720)视频，则每帧的总像素数可以是921600个。一个像素可以包含3种颜色(例如，RGB格式的红色、绿色和蓝色，或YUV格式的亮度和色度)，上述颜色可以分别存储在字节可寻址存储器单元中的连续地址区中。像素的每种颜色的大小是一个字节，并且DS是1。例如，可以定义内核程序以将颜色格式从RGB转换为YUV。视频流的一帧可以被分成多个矩形窗口，并且每个窗口都可以映射线程块。BASE可以是矩形窗口左上角的存储器地址。例如，一帧可以被分成120个160×48个窗口。可以组织线程使得网格和块的尺寸可以分别是8×5×1和160×48×3。BASE可以通过base_frame+blockIdx.x*160+blockIdx.y*1280*48给出。这里，base_frame可以是帧的基址。STRIDE_Y和STRIDE_Z可以分别为1280和1280*720。

在一个实施例中，warp的K个并发线程的第一数据的存储器地址ADDR可以通过使用以下公式中的寻址参数来生成：ADDR＝BASE+threadIdx.x+threadIdx.y*STRIDE_Y+threadIdx.z*STRIDE_Z(当存储器单元112是字可寻址存储器单元时)，以及ADDR＝BASE+(threadIdx.x+threadIdx.y*STRIDE_Y+threadIdx.z*STRIDE_Z)*DS(当存储器单元112是字节可寻址存储器单元时)。该计算可由地址计算单元714执行。向量寄存器可能需要使用前面的列来更新，存储器地址可以不使用向量寄存器而在单个列中生成(例如，一列中的MP)。ThreadIdx.x是线程块x轴上的线程索引，并且是K的倍数，threadIdx.y是线程块y轴上的线程索引，threadIdx.z是线程块z轴上的线程索引。MP可以从序列发生器106接收嵌入在配置中的线程块大小信息(例如，X、Y和XYZ)，并通过以下伪代码计算threadId.x、threadIdx.y、以及threadIdx.z，其中假设K是32。

NUM_EXEC可以根据ceil(XYZ/32)给出，其中ceil(x)是返回大于或等于x的最小整数的上限函数(ceiling function)。在一个实施例中，MP中的计数器740.3可以计算发射的存储器地址的数量，并用作for循环(for-loop)的循环计数器。如上伪代码所示，可以为32个线程的每个warp发射存储器地址ADDR。也就是说，MP可以为线程0生成一个存储器地址(以为线程0到线程31加载数据)，然后为线程32生成另一个存储器地址(以为线程32到线程63加载数据)，等等。需要注意的是，向量操作的粒度可以是32个线程的warp。如果有100个线程，则重复4次向量操作(例如，ceil(100/32))，并且最后一次执行向量操作可能在末尾生成28个字的无效数据。例如，假设存储器单元300是如存储器示例2中的数据大小(datasize，DS)为2的字节可寻址存储器，MP可以向私有存储器地址端口发射4个存储器地址122、186、250、314。使用地址314加载的512位数据可以包含PE可能会生成无效数据的最后一个warp的线程100至线程127的28个字的无效数据。加载额外数据可能不会损害性能，但是存储无效数据可能会损害性能。在一个实施例中，私有存储器地址端口可以使用掩码以避免向存储器存储无效数据和/或从存储器加载无效数据。

图8示意性地示出了根据本公开的实施例的处理器的序列发生器800。序列发生器800可以是图1的序列发生器106的实施例，并且可以解码内核程序、执行解码后的标量指令、将解码后的向量指令打包成配置并将配置传递到列。

序列发生器800可以耦合到任务缓冲器(task buffer)(例如，任务FIFO)816和指令高速缓存(例如，i-Cache)818。内核信息(例如程序的基地址、作业标识符(例如，作业ID)、块标识符(例如，块ID)、以及块索引)可以经由任务缓冲器816传输到序列发生器800。在一个实施例中，任务缓冲器816和指令高速缓存818可以是图1的配置存储器104的一部分。在操作期间，内核信息可以由外部设备写入任务缓冲器816。当任务缓冲器816不为空时，序列发生器800可以在从任务缓冲器816读取内核信息之后开始处理内核，然后从指令高速缓存818检索指令用于解码，并且当程序结束指令被解码时停止。当任务缓冲器816为空时，序列发生器800可以处于空闲。

在一些实施例中，内核信息可以包括一位以指示序列发生器802是否应该以连续模式工作。如果该位被设置，则序列发生器800可以从任务缓冲器816连续读取内核信息并提取内核程序。否则，序列发生器800可以监视列(例如PE阵列114的列)的状态，并且在从任务缓冲器816读取下一内核信息之前等待直到所有列变为未激活。

序列发生器800可包括控制器802、指令缓冲器804、标量指令解码器806、向量指令解码器808、标量处理器810、本地存储器812、以及标量寄存器814。控制器802可以从指令高速缓存818提取指令，并将所提取的指令放入指令缓冲器804中。在一个实施例中，指令缓冲器804可以是循环缓冲器，以保持若干指令(例如，64或另一数量)。在操作期间，例如，控制器802可以在内核的开始处提取64个指令以填充指令缓冲器804。

提取的指令可以包括混合在一起的标量指令和向量指令。指令的某些位(例如，最高有效的4位、最低有效的4位或其他位)可以指定被指定用于执行该指令的硬件单元。控制器802可以检查这些位，并基于所指定的硬件单元确定指令是标量指令还是向量指令。

指令缓冲器804可以具有指令指针，该指令指针指向指令缓冲器804中要处理的下一指令。要处理的下一指令也可以由控制器802中的程序计数器(program counter，PC)820所指向。控制器802可以确定指令是标量指令还是向量指令，并指示标量指令被发送到标量指令解码器806和向量指令被发送到向量指令解码器808。在一些实施例中，标量指令解码器806可以在一个周期内解码一个标量指令，而向量指令解码器808可以在一个周期内解码多个向量指令。例如，在一个实施例中，向量指令解码器808可以在一个周期中解码多达8个向量指令。然而，如果向量指令引用标量寄存器814中的一个或多个寄存器，并且该一个或多个寄存器还没有准备好，则可以插入等待周期。当标量指令解码器806解码一个标量指令时，PC 820可以递增1。当向量指令解码器808解码向量指令时，PC 820可以通过解码后的向量指令的数量递增。

序列发生器800可以按照存储在指令高速缓存818中的次序按顺序地处理指令。由标量指令解码器806解码的标量指令可以由控制器802和标量处理器810即时执行。标量指令可以生成用于配置列以及管理循环和分支的参数。

向量指令解码器808可以解码向量指令以生成用于向量处理单元的配置。向量指令可以配置列中的数据路径，控制数据流并在并行线程中处理数据。例如，用于存储器端口(例如，MP 120)的向量指令可以包括存储器访问指令，例如但不限于LOAD和STORE；用于交换盒(例如，SB 122)的向量指令可以包括数据复制指令，例如但不限于MOVE和FORWARD；用于处理元件(例如，PE 118)的向量指令可以包括算术和逻辑指令，例如但不限于ADD和SUBTRACT等。

图8示出了配置830，配置830作为由向量指令解码器808生成并传送到各列的配置的示例。配置830可以包括用于硬件配置的字段和用于一组立即值的字段。用于硬件配置的字段可以包括要在向量处理单元执行的一个或多个指令。该组立即值可以包括指令的参数。

在一些实施例中，配置可以被广播到所有列。上述列中的一列可以是一组配置的目的列或目标列，每个列可以具有单独的单线，该单线被耦合至序列发生器以用于传输有效位。当选定列的配置缓冲器未满时，标量指令解码器806可以将有效位变为有效以选择特定列。也就是说，当目标列的配置缓冲器(例如，选定列的MP、PE、SB的配置缓冲器)具有可用空间时，可以将目标列的有效位变为有效，并作为配置以被该目标列接收。当有效信号未被有效化时，向量处理单元的配置缓冲器的输入可以被关联到地，以防止向量处理单元中的逻辑切换。

向量处理单元的配置缓冲器大小可以大于一个配置单元大小，以使得不需要将序列发生器800与列进行同步。也就是说，向量处理单元中的每一个可以在任何时间保持多于一个配置，并且向量处理单元的每个列可以相对于序列发生器800以异步方式执行解码后的向量指令。因此，序列发生器800可以在列完成程序执行之前完成配置分派。在至少一个实施例中，序列发生器800可以监视列是激活的还是未激活的，但是不监视列正在执行哪些指令。

可以按照列号的次序一次一个按顺序地选择目的列(destination column)。因为列可以在循环中被链接(例如，从MP 120.1到120.N然后经由垫片存储器116回到MP 120.1链接的MP，从SB 122.1到PE 118.1至SB 122.N到PE 118.N然后经由垫片存储器116回到SB122.1链接的SB和PE)，内核程序的执行可以选择任何列作为起始列。在一个实施例中，可以选择PE阵列114的第一列(例如，MP 120.1、SB 122.1和PE 118.1)来开始内核程序的执行，并且可以按照列号的次序逐个按顺序地选择其他列。

在启动内核之前，可以由外部设备在本地存储器812中设置内核特定参数。本地存储器812可以是随机存取存储器(random access memory，RAM)(例如，DRAM、SRAM等)。当执行一些标量指令时，标量处理器810可以从本地存储器812读取这些参数，处理这些参数，并将这些参数存储在标量寄存器814中。例如，标量处理器810可以执行CONFIGXYZ指令，例如CONFIGXYZ％X、％Y、％XYZ，其中％X、％Y、％XYZ可以是标量寄存器814的三个独立的标量寄存器，该文件被指定用于存储正在处理的内核程序的线程块大小信息。在启动内核之前，X、Y、XYZ值可以存储在本地存储器812中。

标量寄存器814可以由标量处理器810和向量指令解码器808共享。向量指令解码器808可以从标量寄存器814获得内核特定参数，并将该内核特定参数作为配置中的立即数传递给列。例如，大小X、Y、XYZ信息可以作为配置中的立即值传递给列。另外，由执行标量指令的标量处理器810生成的参数(例如，用于配置列)也可以使用标量寄存器814传递到向量指令解码器808。在一些实施例中，标量寄存器814可以包括多个寄存器。例如，在一个实施例中，标量寄存器814可以包括32个16位寄存器，表示为R0到R31。

标量处理器810可以包括标量ALU和加载/存储器单元。在一个实施例中，ALU可以包括整数单元、浮点单元、移动单元和比较单元。这些单元中的每一个都可以在多级流水线中实现。加载/存储器单元也可以实现为多级流水线。加载单元可以从本地存储器812读取数据，并将上述数据存储在标量寄存器814中的标量寄存器。存储器单元可以将标量寄存器814的内容写入本地存储器812。

由标量指令解码器806解码的标量指令可以包括控制流指令，这些指令可以由控制器802执行。在一些实施例中，控制流指令可以包括但不限于重复(repeat)、跳转(jump)、轮询(poll)和屏障(barrier)指令。跳转指令是将执行流程从当前指令序列中的下一指令改变到跳转指令所指向的目的指令。轮询指令是让控制器802停止提取指令并等待直到DMA操作完成(例如，DMA模块102完成)。轮询指令可以将序列发生器800和列进行同步。当执行屏障指令时，控制器802可以停止提取指令并等待直到所有列变为未激活。屏障指令可以将序列发生器800和列进行同步。

在一些实施例中，控制器802可以包括单独的硬件单元，每个硬件单元用于执行分配给控制器802的不同类型的指令。例如，控制器802可以包括用于执行跳转指令的硬件单元、用于执行轮询指令的另一硬件单元、以及用于执行屏障指令的又一硬件单元。

对于重复指令，控制器802可以配备有PC寄存器822、列计数器824、循环计数器826和列号寄存器828。在一个实施例中，这些硬件单元可以包含在被指定用于重复指令的一个硬件单元中。当执行重复指令时，控制器802可以在PC寄存器822中捕获PC 820的当前值，清除列计数器824，将循环计数器826设置为指令要重复执行的总次数，并将列数寄存器828设置为要重复的列数。后两个值可以由重复指令提供。当配置被发送到列时，控制器802可以递增列计数器824。如果列计数器824中的列计数器值等于列号寄存器828中的列号寄存器值，则控制器802可以清除列计数器824并递减循环计数器826。如果循环计数器826中的循环计数器值不为零，则控制器802可以用PC寄存器822中的PC寄存器值替换PC 820中的值，以跳回到循环的起点。如果循环计数器值为零，则控制器802可以退出循环。

图9是根据本公开实施例的处理器使用标量存储器地址以加载或存储向量数据的过程900的流程图。在块902，序列发生器可以向存储器端口发送线程块大小信息和寻址参数(例如，BASE、STRIDE_Y、STRIDE_Z、以及DS)。例如，外部设备可以在序列发生器800的本地存储器812中设置数据的基址以及线程块和网格的大小信息，然后通过在任务缓冲器816中设置线程块索引信息来发起内核程序。序列发生器800可以用于从任务缓冲器816加载线程块索引信息，然后执行内核程序中的标量指令。标量指令可以使序列发生器800从本地存储器812加载数据的基址以及线程块和网格的大小信息；使用从本地存储器812加载的数据的基址以及线程块和网格的大小信息来计算寻址参数；将线程块大小信息和寻址参数存储到标量寄存器中。当用于数据加载过程的加载(LOAD)指令(或用于数据存储过程的存储指令)被解码时，寻址参数可以从标量寄存器捕获，该寻址参数可以与线程块大小信息一起嵌入到存储器端口的配置中，并被传送到存储器端口。

在块904，存储器端口可以基于线程块大小信息和寻址参数(例如，通过地址计算单元714)生成标量存储器地址。该标量存储器地址可以指向warp的第一线程的数据的存储位置。

在块906，存储器接口和存储器单元可以对于数据加载过程，从起始于标量存储器地址的K个连续存储器地址将warp的数据的K个字加载到向量寄存器，或者对于数据存储过程，从向量寄存器将warp的数据的K个字存储到起始于标量存储器地址的K个连续存储器地址。K可以是warp大小，warp大小表示要通过执行向量指令来同时处理的线程数。例如，处理器100的每个向量处理单元可以有向量大小K。也就是说，处理器100的每个向量处理单元可以用于执行warp中的K个并发线程(例如，PE中的K个ALU)。使用标量存储器地址作为起始地址，将用于数据加载过程的K个连续存储器地址的K个数据字可以被加载到向量寄存器，或用于数据存储过程的K个线程的K个数据字可以从向量寄存器存储到K个连续存储器地址。

本公开提供了用于可重构并行处理器(reconfigurable parallel processor，RPP)的装置、系统和方法。例如，RPP的一个实施例可以使用处理元件(PE)阵列作为物理数据路径来处理大量并行数据。物理数据路径可以在每个部分(例如，MP、SB和PE的一列)中相同，从而可以允许内核程序的相依性图表(dependency graph)被映射到虚拟数据路径，该虚拟数据路径可以是物理数据路径的无限重复。另外，内核程序的标量指令可以由序列发生器执行，而无需为执行标量指令的向量处理单元生成任何配置。另外，标量控制流指令也可以由序列发生器执行，而不用浪费任何向量处理单元周期。

RPP的实施例还可以使用垫片存储器来临时存储物理数据路径(例如，处理元件(PE)阵列)的输出，该物理数据路径由虚拟数据路径的一个分段配置，该虚拟数据路径由N个向量执行节点组成。垫片存储器可以用作数据缓冲器(例如，FIFO)，当物理数据路径被虚拟数据路径的下一分段重新配置时，该数据缓冲器将数据反馈回物理数据路径。

RPP的实施例还可以具有存储器单元，该存储器单元具有连接到物理数据路径的每一列(例如，MP、SB和PE)的存储器接口(memory interface，MI)。通过虚拟数据路径访问的所有数据可以存储在存储器单元中。对于虚拟数据路径的每个分段，MP可以被重配置以不同的方式访问存储器单元，而数据可以保持不变。

RPP的各实施例可以被优化以实现单指令多线程(SIMT)处理的大规模并行性。在一个示例中，对于一行32个PE并且每个PE具有32个算术逻辑单元(ALU)，1024个ALU可以被包括在一个RPP核中。在一些实施例中，多核处理器可以包括多个RPP。

RPP的各实施例可以根据重配置机制进行重配置。RPP的各种组件(包括一个或多个重配置计数器)可称为可重配置单元。例如，PE(例如，PE 118)、交换盒(例如，SB 122)和存储器端口(例如，MP 120)中的每一个可以包括一个或多个重配置计数器，例如PE中的计数器206、SB中的计数器520、MP中的计数器740。当线程之间不存在依赖性时，数据处理可以是流水线式的。相同的指令可以被执行多次，直到所有线程(例如，如果线程的总数是1024，32个wrap的线程用于一个可重配置单元)被处理。在各种实施例中，当可重配置单元中的计数器达到编程数(例如NUM_EXEC)时，可重配置单元可以将其配置替换为新的上下文(context)。这种重配置可以在每个PE、SB和MP中以相同的方式进行。可以实现具有最小的切换空闲时间的自重配置(self-reconfiguration)。

因为在已经处理玩所有线程之后仅交换一次配置，所以示例性重配置机制可以减少在配置功耗。这还可以通过在每个可重配置单元的最早时间独立地交换每个可重配置单元来减少配置之间的空闲时间。

在一些实施例中，在共享存储器访问模式下，所有warp可以使用相同的地址加载相同的数据。由于操作的流水线性质，可以只需要执行第一warp的数据加载(LOAD)指令。加载的数据可以与其他warp共享，以减少存储器访问流量和功耗。

各实施例可以提供私有存储器模式以使用标量存储器地址访问向量数据的存储器单元中的顺序存储器空间。因为访问向量数据只需要标量存储器地址，所以私有存储器地址总线可以只需要标量存储器地址的位宽(而不是向量地址总线)，并且私有存储器地址总线可以被称为标量存储器地址总线。另外，私有存储器访问模式下没有存储器争用，因此可以简化存储器接口。另外，私有存储器访问模式不需要MP中的ALU，因此MP也可以被简化(例如，在MP 700中，ALU 718可以仅被实现用于共享存储器访问的地址计算)。另外，因为不需要在前一列中使用MP和/或PE来更新用于地址计算的向量寄存器，所以可以提高对处理器100的列的使用率。

在示例性实施例中，一种方法可以包括：序列发生器向存储器端口发送序列发生器的标量寄存器中存储的线程块大小信息和寻址参数，存储器端口基于线程块大小信息和寻址参数生成标量存储器地址，以及存储器接口和存储器单元从起始于标量存储器地址的K个连续存储器地址将warp的数据的K个字加载到向量寄存器。标量存储器地址可以指向数据warp的第一线程的数据的存储位置。K可以是warp大小，warp大小表示要通过执行向量指令来同时处理的线程数。该数据的大小可以小于或等于字大小。

在一个实施例中，该方法还可以包括：从序列发生器的任务缓冲器加载线程块索引信息，以及在序列发生器执行内核程序中的标量指令以：从序列发生器的本地存储器加载数据的基址以及线程块和网格的大小信息；使用从本地存储器加载的数据的基址以及线程块和网格的大小信息计算寻址参数；将线程块大小信息和寻址参数存储到序列发生器的标量寄存器中。

在一个实施例中，线程块大小信息可以包括x轴上的线程大小的第一值X、y轴上的线程大小的第二值Y、以及X乘以Y和Z的乘法结果XYZ，其中Z是z轴上的线程大小，并且当存储器单元是字可寻址存储器单元时，寻址参数可以包括BASE、STRIDE_Y、以及STRIDE_Z，或者当存储器单元是字节可寻址存储器单元时，寻址参数可以包括BASE、STRIDE_Y、STRIDE_Z、以及DS。BASE可以是内核程序要处理的数据的起始存储器地址，对于字可寻址存储器单元，STRIDE_Y和STRIDE_Z可以分别是y轴和z轴上的存储器地址的间隙，或者对于字节可寻址存储器单元，STRIDE_Y和STRIDE_Z分别是y轴和z轴上的存储器地址的间隙除以DS，并且DS可以是以字节为单位的存储器单元的数据大小。

在一个实施例中，存储器单元可以包括包括多个存储器组，标量存储器地址可以包括作为存储器组索引的组选择位，存储器组索引指向多个存储器组中的存储器组，并且存储器接口可以用于使用从通过组选择位选择的存储器组开始的标量存储器地址，生成多个组地址，其中每个存储器组一个组地址，并且按照相应线程的顺序，将从多个存储器组接收的数据的K个字重新排列。

在一个实施例中，标量存储器地址可以由存储器端口计算，对于字可寻址存储器单元，计算为ADDR＝BASE+threadIdx.x+threadIdx.y*STRIDE_Y+threadIdx.z*STRIDE_Z，或者对于字节可寻址存储器单元，计算为ADDR＝BASE+(threadIdx.x+threadIdx.y*STRIDE_Y+threadIdx.z*STRIDE_Z)*DS，其中threadIdx.x可以是内核程序的线程块中x轴上的线程索引，并且是K的倍数，threadIdx.y可以是线程块中y轴上的线程索引，threadIdx.z可以是线程块中z轴上的线程索引。

在一个实施例中，内核程序的线程块可以被映射到图像的数据集中的矩形窗口，并且BASE是该矩形窗口左上角的存储器地址。

在一个实施例中，存储器端口可以被可以包括标量存储器地址总线的私有存储器地址端口经由存储器接口耦合至存储器单元。标量存储器地址可以经由标量存储器地址总线传送到存储器接口。并且如果存储器单元是字节可寻址存储器单元，则数据大小DS可以经由标量存储器地址总线的数据大小位传送到存储器接口。

在一个实施例中，向量寄存器可以被映射到交换盒的输出数据缓冲器，或是被映射到从所述存储器端口到所述交换盒的读数据连接的虚拟向量寄存器。

在另一示例性实施例中，处理器可以包括序列发生器、被耦合至该序列发生器的存储器端口、以及经由存储器接口耦合至该存储器端口的存储器单元。序列发生器可以用于向存储器端口发送序列发生器的标量寄存器中存储的线程块大小信息和寻址参数。存储器端口可以用于基于线程块大小信息和寻址参数生成标量存储器地址，所述标量存储器地址指向数据warp的第一线程的数据的存储位置。存储器单元和存储器接口可以用于：对于数据加载过程，从起始于标量存储器地址的K个连续存储器地址将warp的数据的K个字加载到向量寄存器，或者对于数据存储过程，从向量寄存器将warp的数据的K个字存储到起始于标量存储器地址的K个连续存储器地址。K可以是warp大小，warp大小表示要通过执行向量指令来同时处理的线程数。该数据的大小可以小于或等于字大小。

在一个实施例中，序列发生器还可以还用于从序列发生器的任务缓冲器加载线程块索引信息，并执行内核程序中的标量指令，以：从序列发生器的本地存储器加载数据的基址以及线程块和网格的大小信息；使用从本地存储器加载的数据的基址以及线程块和网格的大小信息计算寻址参数；以及将线程块大小信息和寻址参数存储到序列发生器的标量寄存器中。

在一个实施例中，线程块大小信息可以包括x轴上的线程大小的第一值X、y轴上的线程大小的第二值Y、以及X乘以Y和Z的乘法结果XYZ，其中Z是z轴上的线程大小，并且当存储器单元是字可寻址存储器单元时，寻址参数可以包括BASE、STRIDE_Y、以及STRIDE_Z，或者当存储器单元是字节可寻址存储器单元时，寻址参数可以包括BASE、STRIDE_Y、STRIDE_Z、以及DS。BASE可以是内核程序要处理的数据的起始存储器地址，对于字可寻址存储器单元，STRIDE_Y和STRIDE_Z可以分别是y轴和z轴上的存储器地址的间隙，或者对于字节可寻址存储器单元，STRIDE_Y和STRIDE_Z分别是y轴和z轴上的存储器地址的间隙除以DS。DS可以是以字节为单位的存储器单元的数据大小。

在一个实施例中，存储器单元可以包括多个存储器组，标量存储器地址可以包括作为存储器组索引的组选择位，存储器组索引指向多个存储器组中的存储器组。存储器接口可以用于使用从通过组选择位选择的存储器组开始的标量存储器地址，生成多个组地址，其中每个存储器组一个组地址。对于数据加载过程，存储器接口还可以用于按照相应线程的顺序，将从多个存储器组接收的数据的K个字重新排列。并且，对于数据存储过程，存储器接口还可以用于将从存储器端口接收的数据的K个字分成多个分组，其中每个存储器组一个分组，每个分组被存储到多个存储器组中的一个存储器组中。

在一个实施例中，标量存储器地址可以由存储器端口计算，对于字可寻址存储器单元，计算为ADDR＝BASE+threadIdx.x+threadIdx.y*STRIDE_Y+threadIdx.z*STRIDE_Z，或者对于字节可寻址存储器单元，计算为ADDR＝BASE+(threadIdx.x+threadIdx.y*STRIDE_Y+threadIdx.z*STRIDE_Z)*DS。threadIdx.x可以是内核程序的线程块中x轴上的线程索引，并且是K的倍数，threadIdx.y可以是线程块中y轴上的线程索引，threadIdx.z可以是线程块中z轴上的线程索引。

在一个实施例中，内核程序的线程块可以被映射到图像的数据集中的矩形窗口，并且BASE是该矩形窗口左上角的存储器地址。

在一个实施例中，存储器端口可以被可以包括标量存储器地址总线的私有存储器地址端口经由存储器接口耦合至存储器单元。标量存储器地址可以经由标量存储器地址总线传送到存储器接口。并且如果存储器单元是字节可寻址存储器单元，则数据大小DS可以经由标量存储器地址总线的数据大小位传送到存储器接口。

在一个实施例中，向量寄存器可以被映射到交换盒的输出数据缓冲器，或是被映射到从所述存储器端口到所述交换盒的读数据连接的虚拟向量寄存器。

在又一个实施例中，一种方法可以包括：序列发生器向存储器端口发送序列发生器的标量寄存器中存储的线程块大小信息和寻址参数，存储器端口基于线程块大小信息和寻址参数生成标量存储器地址，以及存储器接口和存储器单元从向量寄存器将warp的数据的K个字存储到起始于标量存储器地址的K个连续存储器地址。标量存储器地址可以指向数据warp的第一线程的数据的存储位置。K可以是warp大小，warp大小表示要通过执行向量指令来同时处理的线程数。该数据的大小可以小于或等于字大小。

在一个实施例中，该方法还可以包括从序列发生器的任务缓冲器加载线程块索引信息，以及在序列发生器执行内核程序中的标量指令以：从序列发生器的本地存储器加载数据的基址以及线程块和网格的大小信息；使用从本地存储器加载的数据的基址以及线程块和网格的大小信息计算寻址参数；将线程块大小信息和寻址参数存储到序列发生器的标量寄存器中。

在一个实施例中，线程块大小信息可以包括x轴上的线程大小的第一值X、y轴上的线程大小的第二值Y、以及X乘以Y和Z的乘法结果XYZ，Z是z轴上的线程大小，并且当存储器单元是字可寻址存储器单元时，寻址参数可以包括BASE、STRIDE_Y、以及STRIDE_Z，或者当存储器单元是字节可寻址存储器单元时，寻址参数可以包括BASE、STRIDE_Y、STRIDE_Z、以及DS。BASE可以是内核程序要处理的数据的起始存储器地址，对于字可寻址存储器单元，STRIDE_Y和STRIDE_Z可以分别是y轴和z轴上的存储器地址的间隙，或者对于字节可寻址存储器单元，STRIDE_Y和STRIDE_Z分别是y轴和z轴上的存储器地址的间隙除以DS。并且DS可以是以字节为单位的存储器单元的数据大小。

在一个实施例中，存储器单元可以包括多个存储器组，标量存储器地址可以由存储器端口计算，对于字可寻址存储器单元，计算为ADDR＝BASE+threadIdx.x+threadIdx.y*STRIDE_Y+threadIdx.z*STRIDE_Z，或者对于字节可寻址存储器单元，计算为ADDR＝BASE+(threadIdx.x+threadIdx.y*STRIDE_Y+threadIdx.z*STRIDE_Z)*DS。threadIdx.x可以是内核程序的线程块中x轴上的线程索引，并且是K的倍数，threadIdx.y可以是线程块中y轴上的线程索引，threadIdx.z可以是线程块中z轴上的线程索引。

本文描述的技术可以以数字逻辑门中的一个或多个专用集成电路(applicationspecific integrated circuit，ASIC)实现，或者通过执行存储在有形处理器可读存储器存储介质中的指令的处理器实现。在一个实施例中，所公开的方法和操作中的任何一个都可以以包括存储在一个或多个计算机可读存储介质上的计算机可执行指令的软件实现。一个或多个计算机可读存储介质可包括非暂时性计算机可读介质(例如可移动或不可移动磁盘、磁带或盒式磁带、固态驱动器(solid state drive，SSD)、混合硬盘驱动器、CD-ROM、CD-RW、DVD或任何其他有形存储介质)、易失性存储器组件(例如DRAM或SRAM)或非易失性存储器组件(例如硬盘驱动器))。计算机可执行指令可以在处理器(例如，微控制器、微处理器、数字信号处理器等)上执行。此外，本公开的实施例可以用作通用处理器、图形处理器、微控制器、微处理器或数字信号处理器。

应当注意，如本文所使用的，两个组件之间的“耦合”和“连接”(例如一个组件“耦合”或“连接”到另一组件)可以指两个组件之间的电子连接，其可以包括但不限于通过电子布线、通过电子元件(例如，电阻器、晶体管)连接等。

尽管已经在此公开了各方面和实施例，但其他方面和实施例对本领域技术人员将是显而易见的。本文公开的各方面和实施例用于说明而非限制，真正的范围和精神由所附的权利要求书指示。

Claims (20)

1.一种方法，包括：

序列发生器向存储器端口发送所述序列发生器的标量寄存器中存储的线程块大小信息和寻址参数；

所述存储器端口基于所述线程块大小信息和所述寻址参数生成标量存储器地址，所述标量存储器地址指向warp的第一线程的数据的存储位置；以及

存储器接口和存储器单元从起始于所述标量存储器地址的K个连续存储器地址将所述warp的数据的K个字加载到向量寄存器，K是warp大小，所述warp大小表示要通过执行向量指令来同时处理的线程数，其中所述数据的大小小于或等于字大小。

2.根据权利要求1所述的方法，还包括：

从所述序列发生器的任务缓冲器加载线程块索引信息；以及

在所述序列发生器执行内核程序中的标量指令，以：

从所述序列发生器的本地存储器加载数据的基址以及线程块和网格的大小信息，

使用从所述本地存储器加载的所述数据的基址以及所述线程块和网格的大小信息，计算所述寻址参数，以及

将所述线程块大小信息和所述寻址参数存储到所述序列发生器的所述标量寄存器中。

3.根据权利要求2所述的方法，其中，所述线程块大小信息包括x轴上的线程大小的第一值X、y轴上的线程大小的第二值Y、以及X乘以Y和Z的乘法结果XYZ，其中Z是z轴上的线程大小，并且当所述存储器单元是字可寻址存储器单元时，所述寻址参数包括BASE、STRIDE_Y、以及STRIDE_Z，或者当所述存储器单元是字节可寻址存储器单元时，所述寻址参数包括BASE、STRIDE_Y、STRIDE_Z、以及DS，其中BASE是所述内核程序要处理的数据的起始存储器地址，对于所述字可寻址存储器单元，STRIDE_Y和STRIDE_Z分别是y轴和z轴上的存储器地址的间隙，或者对于所述字节可寻址存储器单元，STRIDE_Y和STRIDE_Z分别是y轴和z轴上的存储器地址的间隙除以DS，并且DS是以字节为单位的所述存储器单元的数据大小。

4.根据权利要求3所述的方法，其中，所述存储器单元包括多个存储器组，所述标量存储器地址包括作为存储器组索引的组选择位，所述存储器组索引指向所述多个存储器组中的存储器组，并且其中，所述存储器接口用于使用从通过所述组选择位选择的所述存储器组开始的所述标量存储器地址，生成多个组地址，其中每个存储器组一个组地址，并且按照相应线程的顺序，将从所述多个存储器组接收的数据的K个字重新排列。

5.根据权利要求4所述的方法，其中，所述标量存储器地址由所述存储器端口计算，对于所述字可寻址存储器单元，计算为ADDR＝BASE+threadIdx.x+threadIdx.y*STRIDE_Y+threadIdx.z*STRIDE_Z，或者对于所述字节可寻址存储器单元，计算为ADDR＝BASE+(threadIdx.x+threadIdx.y*STRIDE_Y+threadIdx.z*STRIDE_Z)*DS，其中threadIdx.x是所述内核程序的线程块中x轴上的线程索引，并且是K的倍数，threadIdx.y是所述线程块中y轴上的线程索引，threadIdx.z是所述线程块中z轴上的线程索引。

6.根据权利要求5所述的方法，其中，所述内核程序的所述线程块被映射到图像的数据集中的矩形窗口，并且BASE是所述矩形窗口左上角的存储器地址。

7.根据权利要求1所述的方法，其中，所述存储器端口通过包括标量存储器地址总线的私有存储器地址端口经由所述存储器接口耦合至所述存储器单元，所述标量存储器地址经由所述标量存储器地址总线传送到所述存储器接口，并且如果所述存储器单元是字节可寻址存储器单元，则数据大小DS经由所述标量存储器地址总线的数据大小位传送到所述存储器接口。

8.根据权利要求1所述的方法，其中，所述向量寄存器被映射到交换盒的输出数据缓冲器，或是被映射到从所述存储器端口到所述交换盒的读数据连接的虚拟向量寄存器。

9.一种处理器，包括：

序列发生器；

存储器端口，耦合至所述序列发生器，其中，所述序列发生器用于向所述存储器端口发送所述序列发生器的标量寄存器中存储的线程块大小信息和寻址参数，并且所述存储器端口用于基于所述线程块大小信息和所述寻址参数生成标量存储器地址，所述标量存储器地址指向数据warp的第一线程的数据的存储位置；以及

存储器单元，经由存储器接口耦合至所述存储器端口，并且用于：对于数据加载过程，从起始于所述标量存储器地址的K个连续存储器地址将warp的数据的K个字加载到向量寄存器，或者对于数据存储过程，从所述向量寄存器将warp的数据的K个字存储到起始于所述标量存储器地址的所述K个连续存储器地址，K是warp大小，所述warp大小表示要通过执行向量指令来同时处理的线程数，其中所述数据的大小小于或等于字大小。

10.根据权利要求9所述的处理器，所述序列发生器还用于从所述序列发生器的任务缓冲器加载线程块索引信息，并执行内核程序中的标量指令，以：

从所述序列发生器的本地存储器加载数据的基址以及线程块和网格的大小信息，

使用从所述本地存储器加载的所述数据的基址以及所述线程块和网格的大小信息，计算所述寻址参数，以及

将所述线程块大小信息和所述寻址参数存储到所述序列发生器的所述标量寄存器中。

11.根据权利要求10所述的处理器，所述线程块大小信息包括x轴上的线程大小的第一值X、y轴上的线程大小的第二值Y、以及X乘以Y和Z的乘法结果XYZ，其中Z是z轴上的线程大小，并且当所述存储器单元是字可寻址存储器单元时，所述寻址参数包括BASE、STRIDE_Y、以及STRIDE_Z，或者当所述存储器单元是字节可寻址存储器单元时，所述寻址参数包括BASE、STRIDE_Y、STRIDE_Z、以及DS，其中BASE是所述内核程序要处理的数据的起始存储器地址，对于所述字可寻址存储器单元，STRIDE_Y和STRIDE_Z分别是y轴和z轴上的存储器地址的间隙，或者对于所述字节可寻址存储器单元，STRIDE_Y和STRIDE_Z分别是y轴和z轴上的存储器地址的间隙除以DS，并且DS是以字节为单位的所述存储器单元的数据大小。

12.根据权利要求11所述的处理器，其中，所述存储器单元包括多个存储器组，所述标量存储器地址包括作为存储器组索引的组选择位，所述存储器组索引指向所述多个存储器组中的存储器组，并且其中，所述存储器接口用于使用从通过所述组选择位选择的所述存储器组开始的所述标量存储器地址，生成多个组地址，其中每个存储器组一个组地址，以及

对于数据加载过程，所述存储器接口还用于按照相应线程的顺序，将从所述多个存储器组接收的数据的K个字重新排列，以及

对于数据存储过程，所述存储器接口还用于将从所述存储器端口接收的数据的K个字分成多个分组，其中每个存储器组一个分组，每个分组被存储到所述多个存储器组中的一个存储器组中。

13.根据权利要求12所述的处理器，其中，所述标量存储器地址由所述存储器端口计算，对于所述字可寻址存储器单元，计算为ADDR＝BASE+threadIdx.x+threadIdx.y*STRIDE_Y+threadIdx.z*STRIDE_Z，或者对于所述字节可寻址存储器单元，计算为ADDR＝BASE+(threadIdx.x+threadIdx.y*STRIDE_Y+threadIdx.z*STRIDE_Z)*DS，其中threadIdx.x是所述内核程序的线程块中x轴上的线程索引，并且是K的倍数，threadIdx.y是所述线程块中y轴上的线程索引，threadIdx.z是所述线程块中z轴上的线程索引。

14.根据权利要求13所述的处理器，其中，所述内核程序的所述线程块被映射到图像的数据集中的矩形窗口，并且BASE是所述矩形窗口左上角的存储器地址。

15.根据权利要求9所述的处理器，其中，所述存储器端口通过包括标量存储器地址总线的私有存储器地址端口经由所述存储器接口耦合至所述存储器单元，所述标量存储器地址经由所述标量存储器地址总线传送到所述存储器接口，并且如果所述存储器单元是字节可寻址存储器单元，则数据大小DS经由所述标量存储器地址总线的数据大小位传送到所述存储器接口。

16.根据权利要求9所述的处理器，其中，所述向量寄存器被映射到交换盒的输出数据缓冲器，或是被映射到从所述存储器端口到所述交换盒的读数据连接的虚拟向量寄存器。

17.一种方法，包括：

序列发生器向存储器端口发送所述序列发生器的标量寄存器中存储的线程块大小信息和寻址参数；

所述存储器端口基于所述线程块大小信息和所述寻址参数生成标量存储器地址，所述标量存储器地址指向数据warp的第一线程的数据的存储位置；以及

存储器接口和存储器单元从向量寄存器将warp的数据的K个字存储到起始于所述标量存储器地址的K个连续存储器地址，K是warp大小，所述warp大小表示要通过执行向量指令来同时处理的线程数，其中所述数据的大小小于或等于字大小。

18.根据权利要求17所述的方法，还包括：

从所述序列发生器的任务缓冲器加载线程块索引信息；以及

在所述序列发生器执行内核程序中的标量指令，以：

从所述序列发生器的本地存储器加载数据的基址以及线程块和网格的大小信息，

使用从所述本地存储器加载的所述数据的基址以及所述线程块和网格的大小信息，计算所述寻址参数，以及

将所述线程块大小信息和所述寻址参数存储到所述序列发生器的所述标量寄存器中。

19.根据权利要求18所述的方法，其中，所述线程块大小信息包括x轴上的线程大小的第一值X、y轴上的线程大小的第二值Y、以及X乘以Y和Z的乘法结果XYZ，其中Z是z轴上的线程大小，并且当所述存储器单元是字可寻址存储器单元时，所述寻址参数包括BASE、STRIDE_Y、以及STRIDE_Z，或者当所述存储器单元是字节可寻址存储器单元时，所述寻址参数包括BASE、STRIDE_Y、STRIDE_Z、以及DS，其中BASE是所述内核程序要处理的数据的起始存储器地址，对于所述字可寻址存储器单元，STRIDE_Y和STRIDE_Z分别是y轴和z轴上的存储器地址的间隙，或者对于所述字节可寻址存储器单元，STRIDE_Y和STRIDE_Z分别是y轴和z轴上的存储器地址的间隙除以DS，并且DS是以字节为单位的所述存储器单元的数据大小。

20.根据权利要求19所述的方法，其中，所述存储器单元包括多个存储器组，所述标量存储器地址由所述存储器端口计算，对于所述字可寻址存储器单元，计算为ADDR＝BASE+threadIdx.x+threadIdx.y*STRIDE_Y+threadIdx.z*STRIDE_Z，或者对于所述字节可寻址存储器单元，计算为ADDR＝BASE+(threadIdx.x+threadIdx.y*STRIDE_Y+threadIdx.z*STRIDE_Z)*DS，其中threadIdx.x是所述内核程序的线程块中x轴上的线程索引，并且是K的倍数，threadIdx.y是所述线程块中y轴上的线程索引，threadIdx.z是所述线程块中z轴上的线程索引。