CN117421049A - 具有形成循环数据路径的堆叠的列的可重构并行处理器 - Google Patents

Abstract

提供了用于线程级并行处理的处理器、系统和方法。处理器可以包括以二维列阵列布置的多列向量处理单元，其中，多个列堆叠在第一方向上并排放置，并且每个列堆叠具有在第二方向上堆叠的两个列和临时存储缓冲器。每个列可以包括处理元件(PE)，该PE具有向量算术逻辑单元(ALU)以在并行线程中执行算术运算。在列阵列在第一方向上的第一端，列堆叠中的两个列耦合到临时存储缓冲器以用于单向数据流。在列阵列在第一方向上的第二端，两个列彼此耦合以用于单向数据流。列阵列和临时存储缓冲器可以形成单向循环数据路径。

Description

具有形成循环数据路径的堆叠的列的可重构并行处理器

技术领域

本公开涉及计算机架构，尤其涉及在单指令多线程(single instructionmultiple threads，SIMT)计算系统中使用堆叠的列来形成循环数据路径(circular datapath)的多线程计算机处理器架构。

背景技术

图形处理单元(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)变得更加实用和流行。然而，在许多处理元件和高速缓冲存储器封装在一个处理器中的情况下，难以为处理器中的所有处理元件提供对所有高速缓冲存储器的访问。因此，本领域需要实现考虑了面积、延迟和功耗的布局规划(floorplan)。

发明内容

本公开描述了用于大规模线程级并行处理的处理器的装置、方法和系统。该处理器可能具有大量的专用接口，但缺乏真正的异步分布式总线。进程的向量处理单元可以形成列，并且向量处理单元的列可以访问所有高速缓存块(cache block)，因此可以访问核中的所有存储器。为了优化面积、延迟和功耗，处理器可以具有仔细调整的布局规划，该布局规划使列到列接口(column-to-column interface)创建循环数据流，其中，最后的列经由临时存储缓冲器(temporary storage buffer)环回至第一列。每个列可以具有驱动后续列的输出FIFO的集合。后续列的输入可以驱动列中的多路复用逻辑，该多路复用逻辑可以将这些信号引导到存储器端口、处理元件或引导到输出FIFO和下一列。最后的列可以使其输出FIFO驱动临时存储缓冲器，该临时存储缓冲器然后驱动第一列。

为了优化面积和功耗，列可以被堆叠以形成列堆叠(column stack)，并且列堆叠可以并排放置，并且可以以环形方式布置列，使得输出FIFO到后续列中的多路复用逻辑的路由距离较短。存在需要从一个列驱动到下一列的许多总线。交换逻辑可以在逐位的基础上操作，因此总线可以是位对齐的(bit-aligned)并沿着列边缘散布。

集中式列方法被实现用于处理器。在这种方法中，向量处理单元的列被放置在阵列中，其中，存储器单元被分成多个部分并放置在阵列的两侧。临时存储缓冲器可以在垂直于堆叠列的方向的方向上放置在列阵列的一端。

在示例性实施例中，处理器可以包括以二维列阵列布置的多列向量处理单元，其中，多个列堆叠在第一方向上并排放置，并且每个列堆叠具有在第二方向上堆叠的两个列，以及临时存储缓冲器。每个列可以包括处理元件(processing element，PE)，该PE具有向量算术逻辑单元(arithmetic logic unit，ALU)以在并行线程中执行算术运算。在列阵列在第一方向上的第一端，第一列堆叠可以具有针对第二列进行堆叠的第一列，第一列可以具有耦合到临时存储缓冲器的输出端口，并且第二列可以具有耦合到临时存储缓冲器的输入端口。在列阵列在第一方向上的第二端，第二列堆叠可以具有针对第四列进行堆叠并连接到第四列的第三列，以使数据从第四列流向第三列。对于列阵列中在第一列堆叠和第二列堆叠之间的列堆叠，每个列不连接到在同一列堆叠内针对该列进行堆叠的另一列，而是可以具有连接到第一相邻列的输入端口和连接到第二相邻列的输出端口，第一相邻列和第二相邻列在第一方向上的相对两侧。列阵列和临时存储缓冲器可以形成单向循环数据路径。

附图说明

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

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

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

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

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

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

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

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

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

图8示意性地示出了根据本公开的实施例的处理器的序列发生器(sequencer，也称为程序发生器、分发器、定序器)。

图9示意性地示出了根据本公开的实施例的处理器的布局规划。

图10示意性地示出了根据本公开的实施例的三个连续的列堆叠。

图11示意性地示出了根据本公开的实施例的列堆叠中的两个列到存储器块的连接。

具体实施方式

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

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

图1示意性地示出了根据本公开的实施例的处理器100。处理器100可以包括直接存储器访问(direct memory access，DMA)模块102、配置存储器104、序列发生器106、控制处理器108、程序高速缓存110、存储器单元112、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))。在一个实施例中，物理数据路径配置可以被称为物理数据路径程序，其可以包括物理数据路径中包括的各种组件的单独配置。

尽管未示出，但可以存在将序列发生器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列126.1-126.N的可编程硬件单元或可编程硬件组件。例如，MP120.1、SB 122.1和PE 118.1可以形成PE阵列114的第一列126.1，并且MP 120.N、SB 122.N和PE 118.N可以形成PE阵列114的第N列126.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可以经由相应的MI 124.1-124.N耦合到存储器单元112，以从存储器单元112读取和写入存储器单元112。因此，MI 124.1可以是到PE阵列114的第一列126.1的存储器单元112的网关，以此类推，直到MI 124.N可以是到PE阵列114的第N列126.N的存储器单元112的网关。每个MP 120.1-120.N还可以耦合到相应列中的SB，以向每个列发送数据和从每个列发送数据。例如，MP 120.1可以耦合到SB 122.1，MP 120.2可以耦合到SB 122.2，等等。在一些实施例中，存储器单元112和MI 124.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位。为了支持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的两个输出，并且PE 118.2的另两个输入可以耦合到来自PE 118.1的两个输出，SB 122.3可以示出PE 118.3的两个输入可以耦合到来自MP 120.3的两个输出，并且PE 118.3的另两个输入可以耦合到来自PE118.2的两个输出，依此类推，直到SB 122.N可以示出PE 118.N的两个输入可以耦合到来自MP 120.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的内部存储器中设置。并且序列发生器106可以从内部存储器加载大小，将大小存储在序列发生器106的标量寄存器中，并且将大小作为配置中的立即数传递给列(例如，列中的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位；当M为16时，ALU的输入和输出可以是16位；当M为32时，ALU的输入和输出可以是32位；诸如此类。输入数据块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，并且对于一个数据路径中的所有组件，该NUM_EXEC可以相等。

图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。存储器组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的幂。

每个高速缓存303可以单独地耦合到所有多个MI 124.1-124.N，用于经由连接端口306进行共享存储器访问，并且每个高速缓存304可以单独地耦合到所有多个MI 124.1-124.N，用于经由连接端口308进行私有存储器访问。连接端口306和308中的每一个可以使用两个订阅来标识其连接，其中，第一订阅标识存储器高速缓存(通过存储器高速缓存订阅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和MI124.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是log2(J))。例如，J和L可以分别是八(8)和三(3)，K可以是32，N也可以是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和428可以被配置成使用向量地址。在私有存储器访问模式中，一个向量地址中的K个地址可以根据线程ID以升序连续。因此，在一个实施例中，仅具有最少线程ID的线程的地址可能需要由向量地址指定，并且地址端口404的宽度可以是G位。此外，假设J小于或等于K，则每个组(例如，412和414)的数据端口的宽度可以是(K/J)xM位。由于存储器单元300可以是J路交织存储器，所以地址的最低有效L位可以确定地址的数据可以驻留的存储器组。此处L可以是J减去1的位宽。一个向量数据中的所有K个数据可以均匀地分布在所有存储器组中，并且是可访问的而不具有存储器争用。

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

图4B示意性地示出了根据本公开的实施例的用于私有存储器访问的私有存储器访问接口403。在地址端口404处接收的存储器地址的宽度可以是G位。通过将0、1、2…和J-1添加到地址端口404处的地址，可以创建J个地址。J个地址中的每一个的最低有效L位可以被检查，并且剩余的G减L位(例如，G-L位)可以被分配给其索引匹配最低有效L位的地址A_0到A_J-1中的一个。地址A_0到A_J-1可以经由地址端口410.1到410.J被传递到存储器组。在WData端口406和RData端口408处的向量数据的宽度可以是KxM位。向量中的K个数据可以从0到k-1进行索引。在地址端口404处的地址的最低有效L位可以被添加到K个索引，然后结果的较低L位可以被作为K个数据的索引。在WData端口406处，具有相同索引的数据被级联成(K/J)xM位的数据，并被分配给写数据WD_0到WD_J-1中的一个，该写数据的索引与级联数据的索引相匹配。写数据WD_0到WD_J-1可以经由WData端口412.1到412.J传递到存储器组。在RData端口408处，可以经由RData端口414.1至414.J从每个存储器组接收读数据RD_0至RD_J-1的(K/J)xM位。(K/J)xM位的每个读数据可以被分割成的K/J个M位数据，然后在从KxM位向量创建(K/J)xM位写数据WD_0到WD_J-1的相反过程中使用索引组织成KxM位向量。

图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的地址，拾取其最低有效L位与组索引匹配的所有地址，并通过地址端口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可以是如本文所描述的到PE 118的输入数据，并且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-FIFO 522.7-522.14可以分别映射到向量寄存器IAIB、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)。在一个配置的执行期间，所有计数器可以独立地对通过数据端口的数据数量进行计数。当所有计数器达到配置中指定的执行次数时，可以应用C-FIFO 518中的下一配置。

使用计数器的类似方法可以应用于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可以耦合到SB 122.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示意性地示出了根据本公开的实施例的处理器的存储器端口700。存储器端口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可以被配置成使用基址和线程变量来生成用于私有存储器访问的存储器地址。基址可以是线程的块(或网格)的数据的起始存储器地址。线程变量可以包括块参数，例如但不限于块(或网格)大小。基址和线程变量可以在MP配置中被传递到MP 700。

存储器端口700还可以包括多个计数器740.1-740.8。每个计数器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相关联。

图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等。

在一些实施例中，尽管配置可以具有一个列作为其目的地或目标列，但是配置可以被广播到所有列。每个列可以具有单独的单线，该单线耦合至序列发生器以用于传输有效位。当所选列的配置缓冲器未满时，标量指令解码器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中设置内核特定参数。当执行一些标量指令时，标量处理器810可以从本地存储器812读取这些参数，处理这些参数，并将这些参数存储在标量寄存器814中。标量寄存器814可以由标量处理器810和向量指令解码器808共享。向量指令解码器808可以从标量寄存器814获得内核特定参数，并将该内核特定参数作为配置中的立即数传递给列。此外，由执行标量指令的标量处理器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和列进行同步。

图9示意性地示出了根据本公开的实施例的处理器100的布局规划900。在布局规划900中，处理器100的列126.1至126.N可以分别表示为COL 1至COL 32，其中，N为32。COL 1至COL 32可以形成二维列阵列，其中，多个列堆叠在第一方向上并排放置，并且每个列堆叠具有在第二方向上堆叠的两个列。例如，COL 32可以与COL 1垂直堆叠，列31可以与列2垂直堆叠，以此类推，直到列17可以与列16垂直堆叠，并且列堆叠被水平放置。应该注意，如图9所示，垂直堆叠的列和水平排列的列堆叠只是一个示例。

除了COL 1和COL 32之外，每个列可以通过中间的数据总线耦合到第一方向上两侧的两个相邻列。顶部列COL 17至COL 32可以具有从左到右的数据流，而底部列COL 1至COL 16可以具有从右到左的数据流。数据流的方向可以由虚线箭头指示。因此，COL 1可以是COL 2的前续列，COL 3可以是COL 2的后续列，COL 2可以通过中间的数据总线耦合到COL1，也可以通过中间的数据总线耦合到COL 3。对于最左边的堆叠，底部列COL 16可以通过中间的数据总线直接连接到顶部列COL 17。对于其他15个堆叠，堆叠内的2个列彼此不连接。

最右边的列叠层可以与垫片存储器116接口连接，并且可以保持与除了最左边的堆叠之外的堆叠相同的设计和接口。垫片存储器116可以具有它的耦合到COL 32的输出端口的输入端口，以及耦合到COL 1的输入端口的输出端口。列COL 1至COL 32和垫片存储器116可以形成单向循环数据路径，并且列的编号还示出了列索引可以如何以循环方式递增。

存储器单元112可以被分成两个部分，其中，这两个部分放置在二维列阵列在第二方向上的两侧(例如，上/下)。如图9所示，列阵列顶部的一个部分可以包括在第一方向(例如，左/右)上并排放置的存储器块902.1和901.2，以及列阵列底部的部分可以包括在第一方向上并排放置的存储器块901.3和901.4。存储器块902.1可以包括高速缓存块904.1和904.2，存储器块902.2可以包括高速缓存块904.3和904.4，存储器块902.3可以包括高速缓存块904.5和904.6，并且存储器块902.4可以包括高速缓存块904.7和904.8。对于可以包括8个存储器组的存储器单元112的实施例，例如其中J为8的存储器单元300，每个存储器块可以包含两个存储器组及其对应的存储器高速缓存。例如，存储器块902.1可以包括存储器组302.1和302.2，高速缓存块904.1可以包括用于存储器组302.1的存储器高速缓存303.1和304.1，并且高速缓存块904.2可以包括用于存储器组302.2的存储器高速缓存303.2和304.2。

列堆叠可以通过存储器块侧的908.1到908.32以及列堆叠侧的910.1到910.32的列堆叠到存储器块接口耦合到相邻的存储器块，其中，每个列堆叠在顶部耦合到一个存储器块，在底部耦合到另一个存储器块。例如，COL 32和COL 1的列堆叠可以通过存储器块侧的908.1和列堆叠侧的910.1的列堆叠到存储器块接口耦合到底部的存储器块902.1，同时，COL 32和COL 1的列堆叠可以通过存储器块侧的908.32和列堆叠侧的910.32的列堆叠到存储器块接口耦合到顶部的存储器块902.4。类似地，COL 31和COL 2的列堆叠可以通过接口908.2和910.2耦合到存储器块902.1，并且通过接口908.31和910.31耦合到存储器块920.4；COL 25和COL 8的列堆叠可以通过接口908.8和910.8耦合到存储器块902.1，并且通过接口908.25和910.25耦合到存储器块920.4；COL 24和COL 9的列堆叠可以通过接口908.9和910.9耦合到存储器块902.2，并且通过接口908.24和910.24耦合到存储器块920.3；COL 18和COL 15的列堆叠可以通过接口908.15和910.15耦合到存储器块902.2，并且通过接口908.18和910.18耦合到存储器块920.3；并且COL 17和COL 16的列堆叠可以通过接口908.16和910.16耦合到存储器块902.2，并且通过接口908.17和910.17耦合到存储器块920.3。

两个相邻的存储器块也可以通过一对接口耦合。例如，存储器块902.1可以通过存储器块902.1处的接口906.1和906.2以及存储器块902.2处的接口906.3和906.4耦合到存储器块902.2。此外，存储器块902.3可以通过存储器块902.3处的接口906.5和906.6以及存储器块902.4处的接口906.7和906.8耦合到存储器块902.4。

应当注意，存储器接口可以指MI 124，存储器访问接口可以指私有存储器访问接口403或共享存储器访问接口402，通用术语接口可以指连接的布线(wiring ofconnection)和组件之间的布线连接(wiring connection)的布置。

尽管在布局规划900中未示出，但序列发生器106和处理器100的其他组件可被分组在一起成标量单元块并放置在一侧(例如，顶部、下方、左侧或右侧)。在一个实施例中，序列发生器106可以与垫片存储器116组合并放置在最右侧。在第二方向(例如，水平方向)上，在列堆叠和存储器块之间可以存在路由通道。例如，来自序列发生器106的控制信号和配置可以通过这些水平路由通道路由，并且连接引脚可以放置在列堆叠和标量单元块的顶部(COL 17到COL 32)和底部(COL 1到COL 16)边界上，以访问水平路由通道。对于存储器块，用于控制信号的引脚可以被放置在底部存储器块的顶部，以及顶部存储器块的底部。应当注意，根据与标量单元的距离，可以为一些控制和配置信号线实现零个、一个、或多个流水线延迟级(例如，FIFO缓冲器)或者简单地说，实现延迟级或级。例如，对于COL 1和COL 2，不存在级；对于COL 3至COL 4，可以插入1个级；对于COL 5至COL 7，可以插入2个级；对于COL8至COL 10，可以插入3个级；对于COL 11至COL 13，可以插入4个级；对于COL 14至COL 16，可插入5个级。

图10示意性地示出了根据本公开的实施例的三个连续的列堆叠。三个连续的列堆叠可以包括COL 27和COL 6的一个列堆叠、COL 26和COL 7的另一列堆叠、以及COL 25和COL8的又一列堆叠。每个列可以有两个列到列接口。除了COL 1和COL 32之外，这样的列到列接口中的一个耦合到前续列，而另一个耦合到后续列。COL 1和COL 32中的每一个都可能有一个这样的列到列接口耦合至垫片存储器。每个列到列接口可能有11条总线进入或11条总线离开。

11条总线可以表示为VA、VB、IA、IB、IC、ID、IE、IF、IG、IH和DVA。作为输入VA_IN、VB_IN、IA_IN、IB_IN、IC_IN、ID_IN、IE_IN、IF_IN、IG_IN和IH_IN，前10条总线可直接反馈入列的SB(例如，分别馈入SB 500的输入端口514.1、514.2和524.1-524.8)。这些信号可以被反馈到SB中的功能逻辑中，并且取决于由序列发生器106设置的列配置，被分配到后续列、MP(以及最终的MI和存储器高速缓存)或PE以进行算术运算。最后的总线，dva，可以是用于MP输入的专用地址总线(例如，用于MP 700的输入716的DVA_IN)，其也可以根据配置直接反馈到后续列。作为输出DVA_OUT、VA_OUT、VB_OUT、IA_OUT、IB_OUT、IC_OUT、ID_OUT、IE_OUT、IF_OUT、IG_OUT和IH_OUT，DVA地址总线可以从MP 700的输出722耦合，其他10条总线可以分别从PE 200的输出208.1和208.2以及SB 500的输出526.1到526.8耦合。

应当注意，图10示出了VA、VB、IA、IB、IC、ID、IE、IF、IG、IH和DVA总线，用于总线名称而不是物理位置的说明。在一些实施例中，V*总线(例如，VA和VB)和I*总线(例如，IA、IB、IC、ID、IE、IF、IG和IH)可以是位对齐的并沿着列边缘散布。例如，SB可以包括逐位多路复用逻辑电路，位IA_0[0]、IB_0[0]等可以反馈到相同的功能逻辑中，而位IA_0[1]、IB_0[1]等可以反馈到不同的功能逻辑中，其中，_*指的是总线的向量元素的索引，而[*]指的是向量元素的位索引。因此，这10条总线(例如，两个V*总线和8个I*总线)可以以交织引脚放置模式布置。例如，在图10中，对于上面的列，“VA”实际上并不在“VB”之下。“VA”和“VB”是重叠的，按位排成一行。此外，为了最小化路由，每个列堆叠可以与对齐的引脚物理邻接，以避免列堆叠中的任何额外路由。例如，一个列的IA_IN_0[0]和IA_OUT_0[0]可以在列堆叠的相对侧上具有相同的y坐标。

在具有COL 16和COL 17的最左侧堆叠中，每条总线可能需要垂直路由，以连接堆叠内的下方的列(COL 16)和上方的列(COL 17)。为了平衡每个位之间的垂直距离，不同于镜像方法，对于COL 16和COL 17两者，位和总线顺序可以在相同方向上递增。例如，在COL17中，假设VA输出位从低到高布置为VA_OUT[0]、VA_OUT[1]，并且向上移动直到以VA_OUT[n]结束。因此，对于COL 16，VA输入位可以从顶部到底部布置为VA_IN[n]，VA_IN[n-1]，并且向下移动直到以VA_IN[0]结束。因此，对于n的所有值，COL 16的VA_IN[n]和COL 17的VA_OUT[n]之间的垂直距离可以相同。应当注意，COL 16的VB_IN位可以以相同的位顺序与VA_IN位交织，并且COL 17的VB_OUT位可以以相同的位顺序与VA_OUT位交织。此外，I*总线可以以相同的位对齐和增量位顺序排列，以对于I*总线的所有位具有平衡的垂直距离。

在一些实施例中，DVA总线不与任何其他总线位对齐。然而，对于COL 16和COL 17，DVA总线也可以具有在相同方向上递增的对齐的位。例如，在COL 17中，假设DVA输出位从低到高布置为DVA_OUT[0]、DVA_OUT[1]，并且向上移动直到以DVA_OUT[n]结束。因此，对于COL16，DVA输入位可以从顶部到底部布置为DVA_IN[n]、DVA_IN[n-1]，并且向下移动直到以DVA_IN[0]结束。因此，对于n的所有值，DVA[n]和DVA[n]之间的垂直距离是相同的。

在处理器100中，每个列可以具有对存储器单元112的直接访问，因此，每个列可以具有到顶部和底部的存储器块的直接连接。图10可以示出列堆叠侧的列堆叠到存储器块接口910.1至910.32可以各自包括用于列堆叠中的两个列的两个区段1002.x和1004.x。数字“x”可以与列堆叠到存储器块接口910的标记索引匹配。例如，列堆叠侧的列堆叠到存储器块接口910.6可以包括两个区段1002.6和1004.6，列堆叠侧的列堆叠到存储器块接口910.7可以包括两个区段1002.7和1004.7，列堆叠侧的列堆叠到存储器块接口910.8可以包括两个区段1002.8和1004.8，列堆叠侧的列堆叠到存储器块接口910.25可以包括两个区段1002.25和1004.25，列堆叠侧的列堆叠到存储器块接口910.26可以包括两个区段1002.26和1004.26，列堆叠侧的列堆叠到存储器块接口910.27可以包括两个区段1002.27和1004.27。

列堆叠到存储器块接口910的每个区段可以是用于相应列堆叠中的一个列的接口，以访问顶部的存储器块(例如，存储器块902.1和902.2)或底部(例如，存储器块902.3和902.4)的存储器块。例如，区段1004.6可以用于COL 6访问底部的存储器块，区段1002.6可以用于COL 27访问底部的存储器块，而区段1004.27可以用于COL 6访问顶部的存储器块，区段1002.27可以用于COL 27访问顶部的存储器块。区段1004.7可以用于COL 7访问底部的存储器块，区段1002.7可以用于COL 26访问底部的存储器块，而区段1004.26可以用于COL7访问顶部的存储器块，区段1002.26可以用于COL 26访问顶部的存储器块。区段1004.8可以用于COL 8访问底部的存储器块，区段1002.8可以用于COL 25访问底部的存储器块，而区段1004.25可以用于COL 8访问顶部的存储器块，区段1002.25可以用于COL 25访问顶部的存储器块。

图11示意性地示出了根据本公开的实施例的列堆叠中的两个列到存储器块的布线。所有列经由MI 124.1到124.N(如图3所示)与所有存储器高速缓存接口连接。实施例试图有效地实现总线的物理路由，而不浪费面积和功率。在一些实施例中，从所有32个列到所有8个高速缓存块9041.1-904.8的总线可以通过馈通路由(feedthrough routing)被分布以利用列堆叠和存储器块内的空间。这可以避免面积昂贵的外部路由通道，因为这些通道将不能有效地用逻辑填充。给定集中式的列布局规划，存在两个主要的馈通实现，一个在列堆叠中，一个在存储器块中。如本文所使用的，馈通可以指通过而不与组件的功能逻辑进行电接触。例如，馈通信号线可以在存储器块中具有一个或多个流水线级(例如，FIFO)，但是不具有到存储器块的功能逻辑电路的任何电连接。

在列堆叠中，两个列中的每个可以具有用于顶部高速缓存块的4个接口，以及具有用于底部高速缓存块的另4个接口。例如，如图11所示，在顶部，COL 16可以通过双向布线连接1102.16.5耦合到高速缓存块904.5，通过双向布线连接1104.16.6耦合到高速缓存块904.6，通过双向布线连接1106.16.7耦合到高速缓存块904.7，以及通过双向布线连接1108.16.8耦合到高速缓存块904.8；并且COL 17可以通过双向布线连接1102.17.5耦合到高速缓存块904.5，通过双向布线连接1104.17.6耦合到高速缓存块904.6，通过双向布线连接1106.17.7耦合到高速缓存块904.7，以及通过双向布线连接1108.17.8耦合到高速缓存块904.8。在底部，COL 16可以通过双向布线连接1114.16.1耦合到高速缓存块904.1，通过双向布线连接1116.16.2耦合到高速缓存块904.2，通过双向布线连接1110.16.3耦合到高速缓存块904.3，以及通过双向布线连接1114.16.4耦合到高速缓存块904.4；COL 17可以通过双向布线连接1114.17.1耦合到高速缓存块904.1，通过双向布线连接1116.17.2耦合到高速缓存块904.2，通过双向布线连接1110.17.3耦合到高速缓存块904.3，以及通过双向布线连接1112.17.4耦合到高速缓存块904.4。

如本文所使用的，布线连接可指用于两个组件之间的信号、数据或两者的电布线，且可以包括许多电信号线。布线连接可以通过将一个组件的接口连接到另一个组件的另一接口来实现。为了避免使用许多参考标记扰乱图11，布线连接也可以用来指组件进行布线连接的接口。因此，COL 16可以分别具有用于在底部的高速缓存块904.1、904.2、904.3和904.4的四个接口1114.16.1、1116.16.2、1110.16.3和1112.16.4。此外，COL 16可以分别具有用于在顶部的高速缓存块904.5、904.6、904.7和904.8的四个接口1102.16.5、1104.16.6、1106.16.7和1108.16.8。COL 17可以分别具有用于在底部的高速缓存块904.1、904.2、904.3和904.4的四个接口1114.17.1、1116.17.2、1110.17.3和1112.17.4；并且还可以分别具有用于在顶部的高速缓存块904.5、904.6、904.7和904.8的四个接口1102.17.5、1104.17.6、1106.17.7和1108.17.8。注意，接口的第一索引号可以指列(例如，16对应于COL16)，第二索引号可以指高速缓存块(例如，1对应于904.1，2对应于904.2，等等)。

在每个列堆叠内，底部的列可以通过穿过顶部的列的馈通路由耦合到顶部的高速缓存块，并且顶部的列可以通过穿过底部的列的馈通路由耦合到底部的高速缓存块。例如，因为COL 17在COL 16的顶部，所以COL 16到顶部的高速缓存块的接口可以馈通COL 17(例如，穿过COL 17而没有与COL 17进行任何接触的双向布线连接1102.16.5、1104.16.6、1106.16.7和1108.16.8，在图11中示出为被COL 17遮挡)，并且COL 17到底部的高速缓存块的接口可以馈通COL 16(例如，穿过COL 16而没有与COL 16进行任何接触的双向布线连接1114.17.1、1116.17.2、1110.17.3和1112.17.4，在图11中示出为被COL 16遮挡)。

在存储器块内，可以存在64个接口用于对所有32个列进行布线连接。例如，作为底部部分的一个示例，图11示出了存储器块902.2可以具有高速缓存块904.3的接口1110.16.3和1110.17.3，以分别耦合到COL 16和COL 17；高速缓存块904.4的接口1112.16.4和1112.17.4分别耦合到COL 16和COL 17。此外，存储器块902.1可以具有高速缓存块904.1的接口1114.16.1和1114.17.1，以分别耦合到COL 16和COL 17；高速缓存块904.2的接口1116.16.2和1116.17.2分别耦合到COL 16和COL 17。此外，对于顶部部分，图11示出了存储器块902.3可以具有高速缓存块904.5的接口1102.16.5和1102.17.5，以分别耦合到COL 16和COL 17；高速缓存块904.6的接口1104.16.4和1104.17.6分别耦合到COL16和COL 17。此外，存储器块902.4可以具有高速缓存块904.7的接口1106.16.7和1106.17.7，以分别耦合到COL 16和COL 17；高速缓存块904.8的接口1108.16.8和1108.17.8分别耦合到COL 16和COL 17。

每个存储器块还可以为左侧的列实现馈通路由以访问右侧的存储器块，并且为右侧的列实现馈通路由以访问左侧的存储器块。参考回图9，在布局规划900中，中心可以在顶部的存储器块902.3和902.4之间，在中间的COL 24和COL 9的列堆叠和COL 25和COL 8的列堆叠之间，以及在底部的存储器块902.1和902.2之间。因此，布局规划900示出COL 9至COL24和存储器块902.2和902.3可以位于左侧，COL 1至COL 8和COL 25至COL32和存储器块902.1和902.4可以位于右侧。例如，对于存储器块中的馈通路由，在图11中，用于布线连接1110.16.3、1112.16.4、1110.17.3和1112.17.4的实线可以指示这些布线连接可以耦合到存储器块902.2内的接口，而用于布线连接1114.16.1、1114.17.1、1116.16.2和1116.17.2的虚线可以指示这些布线连接可以馈通存储器块902.2并耦合到存储器块902.1内的接口。此外，用于布线连接1102.16.5、1102.17.5、1104.16.6和1104.17.6的实线可以指示这些布线连接可以耦合到存储器块902.3内的接口，而用于布线连接1106.16.7、1106.17.7、1108.16.8和1108.17.8的虚线可以指示这些布线连接可以馈通存储器块902.3并耦合到存储器块902.4内的接口。

应当注意，在一些实施例中，存储器块可以在一个功能逻辑块中实现(例如，在硬件描述语言中，例如VERILOG)，并且列堆叠也可以在一个功能逻辑块中实现。这些功能逻辑块之间的连接可以通过在功能逻辑块的边界处放置引脚并将这些引脚连接为电连接来实现。参考回图10，列堆叠侧的列堆叠到存储器块接口910的区段1002和1004以及对应的接口908可以包括用于布线连接1102、1104、1106、1108、1110、1112、1114和1116连接的引脚以连接通过。例如，布线连接1102.16.5、1104.16.6、1106.16.7和1108.16.8可以通过接口910.17的区段1004.17的引脚和908.17的对应引脚来连接；布线连接1102.17.5、1104.17.6、1106.17.7和1108.17.8可以通过接口910.17的区段1002.17的引脚和对应的908.17的引脚连接；布线连接1114.16.1、1116.16.2、1110.16.3、1112.16.4可以通过接口910.16的区段1004.16的引脚和对应的908.16的引脚连接；布线连接1114.17.1、1116.17.2、1110.17.3、1112.17.4可以通过接口910.16的区段1002.16的引脚和对应的908.16的引脚连接。

此外，在相邻存储器块之间的边界处可以存在引脚，例如，接口906.1和906.3可以包括用于布线连接1114.16.1和1114.17.1的引脚以连接通过，接口906.2和906.4可以包括用于布线连接1114.16.2和1114.17.2的引脚以连接通过，接口906.5和906.7可以包括用于布线连接1106.16.7和1106.17.7的引脚以连接通过，以及接口906.6和906.8可以包括用于布线连接1108.16.8和1108.17.8的引脚以连接通过。

应当注意，耦合到所有8个高速缓存块的一个列的布线连接可以包括用于共享存储器访问接口402的数据端口422.1-422.J、424.1-424.J和地址端口426.1-426.J的连接，以及用于私有存储器访问接口403的数据端口412.1-412.J、414.1-414.J和地址端口410.1-410.J的连接。

在一些实施例中，由于列堆叠和存储器块的大小，所有信号可能很难在单个周期内穿越接口。在至少一个实施例中，延迟级可以在布线连接中被放置在列边界和高速缓存块的边界(例如，高速缓存块中的寄存器的边界)之间。例如，布线连接1114.16.1、1116.16.2、1110.16.3、1112.16.4、1114.17.1、1116.17.2、1110.17.3和1112.17.4中的每一个可以在存储器块902.2内具有三个延迟级(例如，FIFO缓冲器)，并且布线连接1102.16.5、1104.16.6、1106.16.7、1108.16.8、1102.17.5、1104.17.6、1106.17.6和1108.17.8中的每一个可以在存储器块902.3内具有三个延迟级。

在示例性实施例中，处理器可以包括以二维列阵列布置的多列向量处理单元，其中，多个列堆叠在第一方向上并排放置，并且每个列堆叠具有在第二方向上堆叠的两个列，以及临时存储缓冲器。每个列可以包括处理元件(processing element，PE)，该PE具有向量算术逻辑单元(arithmetic logic unit，ALU)以在并行线程中执行算术运算。在列阵列在第一方向上的第一端，第一列堆叠可以具有针对第二列进行堆叠的第一列，第一列可以具有耦合到临时存储缓冲器的输出端口，并且第二列可以具有耦合到临时存储缓冲器的输入端口。在列阵列在第一方向上的第二端，第二列堆叠可以具有针对第四列进行堆叠并连接到第四列的第三列，以使数据从第四列流向第三列。对于列阵列中在第一列堆叠和第二列堆叠之间的列堆叠，每个列不连接到在同一列堆叠内针对该列进行堆叠的另一列，而是可以具有连接到第一相邻列的输入端口和连接到第二相邻列的输出端口，第一相邻列和第二相邻列在第一方向上的相对两侧。列阵列和临时存储缓冲器可以形成单向循环数据路径。

在一个实施例中，每个列还可以包括提供向量存储器操作的存储器端口(memoryport，MP)和为PE提供向量数据交换的交换盒(switch box，SB)。

在一个实施例中，处理器还可以包括被分成两个部分的存储器单元，这两个部分可以放置在列阵列在第二方向上的两个相对侧上。

在一个实施例中，存储器单元的每个部分可以包含在第一方向上并排放置的两个存储器块。

在一个实施例中，存储器单元可以包括分布在存储器块中的多个存储器组，并且每个列可以具有耦合到所有存储器组的存储器接口。

在一个实施例中，在每个列堆叠内，用于一个列访问列阵列的另一侧上的存储器组的布线连接可以穿过列堆叠内的另一列以到达列堆叠的另一侧上的接触引脚。

在一个实施例中，每个列堆叠被夹在两个存储器块之间，两个存储器块在两侧与列堆叠相邻，并且每个存储器块可以包含用于列堆叠中不在该存储器块旁边的列的布线连接，以穿过该存储器块到达在该存储器块旁边的另一存储器块。

在一个实施例中，每个存储器组可以具有相关联的存储器高速缓存，用于每个列的布线连接可以连接到所有存储器高速缓存。

在一个实施例中，每个存储器高速缓存可以具有用于共享存储器访问的第一高速缓存和用于私有存储器访问的第二高速缓存，并且每个存储器接口可以包含共享存储器访问接口和私有存储器访问接口，并且用于每个列的布线连接可以包括在第一高速缓存和共享存储器访问接口之间的布线连接，以及在第二高速缓存和私有存储器访问接口之间的布线连接。

在一个实施例中，布线连接可以包括每个存储器块中的一个或多个延迟级。

在一个实施例中，多个列中的一个列的列到列接口可以包括用于该列中的MP的地址总线、用于来自该列中的PE的输出的第一类型数据总线、以及用于SB到SB数据路由的第二类型数据总线。

在一个实施例中，第一类型的数据总线和第二类型的数据总线可以位对齐，其中，不同数据总线的相同位置处的位被分组以反馈到同一功能逻辑。

在一个实施例中，用于列到列接口的第一类型数据总线和第二类型数据总线的引脚可以以交织布置模式(interleaving placement pattern)布置。

在一个实施例中，除了第二列堆叠中的第三列和第四列之外，用于列的相对侧上的第一类型数据总线和第二类型数据总线的引脚可以被对齐，以避免列堆叠中的任何额外路由。

在一个实施例中，对于第二列堆叠中的第三列和第四列，不同于镜像布置，用于进入第四列的数据流的列到列接口的引脚用于第三列出来的数据流的列到列接口的引脚的相同方向上可以被对齐递增，以平衡距离。

在一个实施例中，处理器还可以包括序列发生器，该序列发生器用于解码包括标量指令和向量指令的指令，执行解码后的标量指令，并将解码后的向量指令打包为配置，并将配置发送到列阵列中的列。

在一个实施例中，处理器还可以包括一对路由通道，每个路由通道放置在列阵列和存储器单元在列阵列在第二方向上的任一侧的一部分之间。

在一个实施例中，用于列的配置可以经由该对路由通道传输。

在一个实施例中，用于该对路由通道中的连接的引脚可以被放置在列堆叠面向该对路由通道中的任一路由通道的边界上。

在一个实施例中，取决于列到序列发生器的距离，该对路由通道中的连接可以包括零个、一个、或多个延迟级。

本文描述的技术可以以数字逻辑门中的一个或多个专用集成电路(applicationspecific integrated circuit，ASIC)实现，或者通过执行存储在有形处理器可读存储器存储介质中的指令的处理器实现。

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

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

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

Claims (20)

1.一种处理器，包括：

多列向量处理单元，以二维列阵列布置，其中，多个列堆叠在第一方向上并排放置并且每个列堆叠具有在第二方向上堆叠的两个列，每个列包括处理元件(PE)，所述处理元件具有向量算术逻辑单元(ALU)以在并行线程中执行算术运算；以及

临时存储缓冲器，其中，在所述列阵列在所述第一方向上的第一端，第一列堆叠具有针对第二列进行堆叠的第一列，所述第一列具有耦合到所述临时存储缓冲器的输出端口，并且所述第二列具有耦合到所述临时存储缓冲器的输入端口，

在所述列阵列在所述第一方向上的第二端，第二列堆叠具有针对第四列进行堆叠并连接到所述第四列的第三列，以使数据从所述第四列流向所述第三列，

对于所述列阵列中在所述第一列堆叠和所述第二列堆叠之间的列堆叠，每个列不连接到在同一列堆叠内针对所述列进行堆叠的另一列，而是具有连接到第一相邻列的输入端口和连接到第二相邻列的输出端口，所述第一相邻列和所述第二相邻列在所述第一方向上的相对两侧，以及

所述列阵列和所述临时存储缓冲器形成单向循环数据路径。

2.根据权利要求1所述的处理器，其中，每个列还包括提供向量存储器操作的存储器端口(MP)和为所述PE提供向量数据交换的交换盒(SB)。

3.根据权利要求2所述的处理器，还包括被分成两个部分的存储器单元，所述两个部分被放置在所述列阵列在所述第二方向上的两个相对侧上。

4.根据权利要求3所述的处理器，其中，所述存储器单元的每个部分包含在所述第一方向上并排放置的两个存储器块。

5.根据权利要求4所述的处理器，其中，所述存储器单元包括分布在所述存储器块中的多个存储器组，并且每个列具有耦合到所有存储器组的存储器接口。

6.根据权利要求5所述的处理器，其中，在每个列堆叠内，用于一个列访问所述列阵列的另一侧上的存储器组的布线连接穿过所述列堆叠内的另一列以到达所述列堆叠的另一侧上的接触引脚。

7.根据权利要求5所述的处理器，其中，每个列堆叠被夹在两个存储器块之间，所述两个存储器块在两侧与所述列堆叠相邻，并且每个存储器块包含用于列堆叠中不在所述存储器块旁边的列的布线连接，以穿过所述存储器块到达在所述存储器块旁边的另一存储器块。

8.根据权利要求5所述的处理器，其中，每个存储器组具有相关联的存储器高速缓存，用于每个列的布线连接被连接到所有存储器高速缓存。

9.根据权利要求8所述的处理器，其中，每个存储器高速缓存具有用于共享存储器访问的第一高速缓存和用于私有存储器访问的第二高速缓存，并且每个存储器接口包含共享存储器访问接口和私有存储器访问接口，并且用于每个列的所述布线连接包括在所述第一高速缓存和所述共享存储器访问接口之间的布线连接，以及在所述第二高速缓存和所述私有存储器访问接口之间的布线连接。

10.根据权利要求8所述的处理器，其中，所述布线连接包括每个存储器块中的一个或多个延迟级。

11.根据权利要求2所述的处理器，其中，所述多个列中的一个列的列到列接口包括用于所述列中的所述MP的地址总线、用于来自所述列中的所述PE的输出的第一类型数据总线、以及用于SB到SB数据路由的第二类型数据总线。

12.根据权利要求11所述的处理器，其中，所述第一类型数据总线与所述第二类型数据总线位对齐，其中，不同数据总线的相同位置处的位被分组以被反馈到同一功能逻辑。

13.根据权利要求12所述的处理器，其中，用于所述列到列接口的所述第一类型数据总线和所述第二类型数据总线的引脚以交织放置模式布置。

14.根据权利要求12所述的处理器，其中，除了所述第二列堆叠中的所述第三列和所述第四列之外，用于列的相对侧上的所述第一类型数据总线和所述第二类型数据总线的引脚被对齐以避免列堆叠中的任何额外路由。

15.根据权利要求12所述的处理器，对于所述第二列堆叠中的所述第三列和第四列，不同于镜像布置，用于流入所述第四列的数据的列到列接口的引脚与用于流出所述第三列的数据的列到列接口的引脚在相同方向上递增后对齐，以平衡距离。

16.根据权利要求1所述的处理器，还包括序列发生器，所述序列发生器用于：

解码指令，所述指令包括标量指令和向量指令；

执行解码后的标量指令；以及

将解码后的向量指令打包为配置，并将所述配置发送到所述列阵列中的列。

17.根据权利要求16所述的处理器，还包括一对路由通道，每个路由通道放置在所述列阵列和所述存储器单元在所述列阵列在所述第二方向上的任一侧的一部分之间。

18.根据权利要求17所述的处理器，其中，所述列的所述配置经由所述一对路由通道传输。

19.根据权利要求18所述的处理器，其中，用于所述一对路由通道中的连接的引脚被放置在所述列堆叠面向所述一对路由通道中的任一路由通道的边界上。

20.根据权利要求17所述的处理器，其中，取决于列到所述序列发生器的距离，所述一对路由通道中的连接包括零个、一个、或多个延迟级。