CN107229463B - 计算设备和相应计算方法 - Google Patents

Abstract

本发明提供一种计算设备和相关计算方法。计算设备用于执行并行计算，包含一组线程处理单元以及存储器混移引擎，存储器混移引擎包含储存数据元素阵列的多个寄存器以及输入选择器阵列，每一输入选择器耦接于多个寄存器的对应子集，并经由输出线耦接于对应线程处理单元，其中，依据第一控制信号，每一输入选择器自多个寄存器的对应子集发送至少第一数据元素至对应线程处理单元，以及依据第二控制信号，每一输入选择器自多个寄存器的另一子集发送至少第二数据元素至对应线程处理单元，其中另一子集经由其他多个输入线耦接于另一输入选择器。本发明的计算设备和相关计算方法可以实现高效的数据访问。

Description

计算设备和相应计算方法

【交叉引用】

本申请要求申请日为2016年3月24日，美国临时申请号为62/312,567的美国临时申请案的优先权，上述临时申请案的内容一并并入本申请。

【技术领域】

本发明的实施例有关于并行计算设备(parallel computing device)和该并行计算设备所执行的计算方法。

【背景技术】

并行计算已经广泛实现于现代计算系统中。为支持高效的并行计算，多个并行编程模型被开发出来，用于程序员编写跨异构平台执行的代码；例如，这样的平台可包含中央处理单元(CPU)、图形处理单元(GPU)、数字信号处理器(DSP)、现场可编程门阵列(FPGA)、硬件加速器等的组合。常用并行编程模型包含开放计算语言(Open Computing Language，简写为OpenCLTM)、OpenCL的变体和扩展等。

通常，并行编程模型建立在并行计算平台模型的基础上。一个并行计算平台模型，例如采用OpenCL的模型，包含耦接于一组计算资源(computational resource)的主机。计算资源进一步包含一组计算设备，且每一计算设备包含一组计算单元。每一计算单元进一步包含一组处理元件。通常，主机执行串行代码，并发布命令至计算设备。计算设备响应命令执行并行代码，也称为内核。内核是程序中声明的功能，且可由多个处理元件在多个进程中执行。相同的内核可以作为同一工作组(workgroup)的多个工作项执行。同一工作组的多个工作项共享本地存储器中的数据，并通过工作组障碍(barrier)彼此同步。

上述并行编程模型是程序员写出高效的并行代码的有力工具。然而，传统的硬件平台不适合具有特定数据访问模式的某些内核，并且对于工作组的分配不灵活。因此，需要提高并行计算系统的效率。

【发明内容】

依据本发明的示范性实施例，提出一种计算设备和相应计算方法以解决上述问题。

依据本发明的一个实施例，提出一种计算设备，用于执行多个并行计算，计算设备包含一组线程处理单元；以及存储器混移引擎，耦接于所述组线程处理单元，所述存储器混移引擎包含：多个寄存器组成的寄存器阵列，储存自存储器缓冲器获得的多个数据元素的阵列；以及多个输入选择器组成的输入选择器阵列，每一输入选择器经由多个输入线耦接于所述多个寄存器的对应子集，并经由一个或多个输出线耦接于一个或多个对应线程处理单元，其中，依据第一控制信号，每一输入选择器自所述多个寄存器的所述对应子集发送至少第一数据元素至所述一个或多个对应线程处理单元，以及依据第二控制信号，每一输入选择器自所述多个寄存器的另一子集发送至少第二数据元素至所述一个或多个对应线程处理单元，其中所述另一子集经由其他多个输入线耦接于另一输入选择器。

依据本发明的另一实施例，提出一种计算设备，用于执行多个并行计算，计算设备包含控制单元，分配多个工作组至一组批处理；以及所述组批处理耦接于所述控制单元，每一批处理包含：程序计数器，由分配给所述批处理的M个工作组共享，其中M是依据可配置批处理设置决定的正整数；一组线程处理单元，并行执行所述M个工作组中的每一个工作组中的多个工作项的子集；以及溢出存储器，当所述M个工作组中的一个或多个工作组遇到同步障碍时，储存所述M个工作组的中间数据。

依据本发明的另一实施例，提出一种计算方法，用于利用多个线程处理单元执行多个并行计算，计算方法包含：接收将多个数据元素组成的数据元素阵列识别为输入的指令；响应所述指令，自存储器缓冲器加载所述数据元素阵列至多个寄存器组成的寄存器阵列，其中所述寄存器阵列耦接于多个输入选择器组成的输入选择器阵列，且每一输入选择器经由多个输入线耦接于所述多个寄存器的对应子集；依据第一控制信号，每一输入选择器自所述多个寄存器的所述对应子集发送至少第一数据元素至一个或多个对应线程处理单元；以及依据第二控制信号，所述每一输入选择器自所述多个寄存器的另一子集发送至少第二数据元素至所述一个或多个对应线程处理单元，其中所述另一子集经由其他多个输入线耦接于另一输入选择器。

依据本发明的另一实施例，提出一种计算方法，用于利用多个线程处理单元执行多个并行计算，计算方法包含：分配多个工作组至多个批处理组成的一组批处理，其中所述多个批处理中的至少一个被分配给共享程序计数器的M个工作组，其中M是依据可配置批处理设置决定的正整数；通过一组线程处理单元并行执行所述M个工作组中的每一个工作组中的多个工作项的子集；以及响应所述M个工作组中的一个或多个工作组遇到同步障碍的发现，将所述M个工作组的中间数据储存至溢出存储器中，并将所述M个工作组中的多个工作项的下一子集加载至所述多个线程处理单元中用于并行执行。

本发明的计算设备和相应计算方法可以实现高效的数据访问。

【附图说明】

图1示出根据一个实施例的耦合到主机和系统存储器的计算设备的示意图。

图2示出了根据一个实施例的计算设备的进一步细节。

图3示出了根据一个实施例的MSE的示例。

图4示出了根据一个实施例的MSE的另一示例。

图5A和图5B根据一个实施例结合矩阵乘法示出了MSE的数据传输模式的示例。

图6是示出根据一个实施例的由计算设备执行的方法的流程图。

图7示出了根据一个实施例的具有可配置数量的工作组分配的计算单元。

图8A和图8B示出根据一个实施例的工作组分配到批处理中的示例。

图9示出根据一个实施例的由具有用于并行计算的可配置工作组分配的计算设备执行的方法的流程图。

【具体实施方式】

在说明书及权利要求书当中使用了某些词汇来指称特定的组件。所属领域中的技术人员应可理解，制造商可能会用不同的名词来称呼同样的组件。本说明书及权利要求书并不以名称的差异来作为区分组件的方式，而是以组件在功能上的差异来作为区分的基准。在通篇说明书及权利要求书当中所提及的「包含」是开放式的用语，故应解释成「包含但不限定于」。另外，「耦接」一词在此包含任何直接及间接的电气连接手段。因此，若文中描述第一装置耦接于第二装置，则代表第一装置可直接电气连接于第二装置，或透过其它装置或连接手段间接地电气连接至第二装置。

在信号和图像处理中以及在科学和工程的各个领域中经常出现的计算问题涉及滤波器的操作。例如，滤波器可以是移动平均(moving average)滤波器、加权移动平均滤波器、有限脉冲响应(finite impulse response，简写为FIR)滤波器等。滤波器可以具有K个符号(tap)，其中K是大于1的整数，并且滤波器可以是应用于数据系列中的每个数据点(data point)。并行计算设备可以将滤波操作划分为多个任务或线程，其中每个计算单元处理一个线程的计算。因此，计算单元在下文中被称为线程处理单元。计算的一个特征是大多数数据点在计算中重复使用。在常规系统中，具有N个线程处理单元的并行计算设备执行用于访问其源操作数(operand)的每通道(per-lane)存储器访问；例如，每个线程处理单元可以从用于执行K个符号滤波操作的存储器单独访问其K个源操作数。这种“每通道”访问使得从存储器重复访问相同的数据点。对于相同的数据的重复存储器访问不仅在速度方面，而且在功耗方面同样是低效的。

低效的数据访问问题通常不仅发生在数据过滤中，而且在对象发现、一般的图像处理、矩阵乘法以及大范围的其他计算中都可以看到。在下文中，使用数据过滤作为示例来简化描述和说明；然而，应当理解，本发明的实施例可应用于涉及通过并行操作的不同线程处理单元重复使用相同数据元素的广泛范围的计算。

本发明的实施例提供了一种用于并行计算系统中的高效工作执行的系统和方法。以下描述包括两个部分。第一部分描述了一种存储器混移引擎(memory shuffle engine,简写为MSE)，其对于多个线程处理单元重复使用的数据元素实现高效的数据访问。第二部分描述了用于设置由并行执行的可配置数量的工作组共享的程序计数器批处理的机制。

图1示出根据一个实施例的耦合到主机150和系统存储器170的计算设备100的示意图。计算设备100的图示已经被简化；但是应当理解，为了便于说明，计算设备100可以包括从图1中省略的更多的组件。计算设备100的示例包括但不限于图形处理单元(GPU)、数字信号处理器(DSP)、图像处理器、专用应用指令集处理器(application-specificinstruction set processor，简写为ASIP)等。DSP和ASIP执行信号、图像和/或多媒体处理操作。DSP和ASIP都可以编程。ASIP的示例是执行由系统支持的专门功能的专用硬件加速器；例如，编码和解码。GPU执行图形处理任务；例如，创建3D场景的2D光栅表示。图形处理可以被称为3D图形流水线或渲染流水线。3D图形流水线可以通过为了加速计算而定制的固定功能硬件和通用可编程硬件的组合来实现，以允许图形渲染中的灵活性。通用可编程硬件也称为着色器硬件(shader hardware)。除了渲染图形之外，着色器硬件还可以执行通用计算任务。

计算设备100包括用于执行单指令-多数据(single-instruction-multiple-data，简写为SIMD)和/或单指令-多线程(single-instruction-multiple-thread，简写为SIMT)操作的并行执行硬件。在一个实施例中，计算设备100包括一个或多个计算单元110。每个计算单元110还包括一组N个线程处理单元115，诸如算术逻辑单元(Arithmetic LogicUnit，简写为ALU)，用于对具有多个数据集(例如，N个数据集)或线程的相同指令执行并行计算。每个计算单元110还包括本地存储器116，用于存储在相同工作组的工作项之间共享的数据。在一个实施例中，N个线程处理单元115可以是来自同一工作组的分配任务；或者，N个线程处理单元115可以是来自多于一个工作组的分配任务。

在一个实施例中，计算设备100还包括控制器120，也称为存储器混移控制器。控制器120的一个职责是控制从存储器缓冲器140到线程处理单元115的数据加载。存储器缓冲器140可以位于本地存储器116中；或者，存储器缓冲器140可以位于系统存储器170中，如图1中虚线框所示。系统存储器170可以是动态随机存取存储器(dynamic random accessmemory，简写为DRAM)或其他易失性或非易失性存储器，并且通常是芯片外的；即，在与计算设备100不同的芯片上。相反，本地存储器116在芯片上；也就是说，它与计算设备100在相同的芯片上。本地存储器116的示例是静态随机存取存储器(static random access memory，简写为SRAM)。其他易失性或非易失性存储器也可以用作本地存储器116。

在一个实施例中，计算设备100还包括存储器混移引擎(MSE)180。MSE180支持涉及使用数据阵列用于并行计算的大范围操作。MSE 180用作存储器缓冲器140和线程处理单元115之间的中间设备。响应于来自控制器120的控制信号，MSE180可从存储器缓冲器140加载数据阵列，并将数据阵列的适当数据元素发送到每个线程处理单元115。关于MSE 180的进一步细节将提供如下。

在一个实施例中，主机150可以包括一个或多个中央处理单元(CPU)。主机150可以向计算设备100发出命令以指示计算设备100执行并行计算。在一些实施例中，计算设备100和主机150可以集成到芯片上系统(system-on-a-chip，简写为SoC)平台中。在一个实施例中，SoC平台可以是移动计算和/或通信设备(例如，智能电话、平板电脑、膝上型计算机、游戏设备等)、桌面计算系统、服务器计算系统或云计算系统的一部分。

图2示出了根据一个实施例的计算设备100的进一步细节。在该实施例中，示出了计算设备中的一个计算单元110。应当理解，计算设备100可以包括任何数量的计算单元110。指令解码器210解码指令以供线程处理单元115执行。控制器120从存储器缓冲器140加载指令的源操作数到线程处理单元115中，其中存储器缓冲器140可以是芯片外(例如，在系统存储器170中)或在芯片上(例如，如图2所示的本地存储器116中)。一些指令可以在寄存器文件230中具有源操作数(一些指令的源操作数可以放在寄存器文件230中)。在一个实施例中，响应于导致数据阵列访问(例如，滤波操作指令)的指令，控制器120使用两个步骤加载源操作数：首先，将源操作数从存储器缓冲器140加载到MSE180中，然后将源操作数从MSE180加载到线程处理单元115中。因为使用MSE180去除了对存储器缓冲器140中的相同数据元素的重复访问，经由MSE 180的存储器访问比每通道存储器访问更高效。

在一个实施例中，MSE 180在控制器120的控制下把将在滤波操作中使用的所有数据元素从存储器缓冲器140加载到MSE 180中。例如，MSE 180可被用于3个符号(3-tap)滤波操作，其被公式化为：filter_output(i)＝(d(i-1)+d(i)+d(i+1))/3,其中i＝1,并且N是被分配以并行执行滤波操作的线程处理单元115的数量。MSE 180将所有(N+2)数据元素d(0)、d(1)、...、d(N+1)加载到其内部触发器(即MSE寄存器阵列282(也称为“寄存器”))。当线程处理单元115准备好执行滤波操作时，MSE 180在三个时刻将三个连续的数据元素作为源操作数输出到每个线程处理单元115。

在一个实施例中，对于K个符号滤波操作，MSE 180可将将在滤波操作中使用的所有数据元素加载到其内部寄存器阵列282中。尽管K可以是任何正整数，但实际上，K的数量受到MSE 180中的寄存器282的数量以及寄存器阵列282和输入选择器283之间的连接数量的限制。将参考图3和图4提供输入选择器的更多细节。

图3示出了根据一个实施例的MSE 380的示例。MSE 380是图1和图2的MSE 180的示例。在该示例中的MSE 380支持K个符号滤波器的操作，其中K＝3。应当理解，K＝3被用于简化描述和说明的非限制性示例。还应当理解，MSE 380可以包括未示出以简化图示的附加电路。

在该示例中，MSE 380包括用于存储(N+2)个数据元素的(N+2)个MSE寄存器282和用于将数据元素输出到相应的N个线程处理单元115的一组N个复用器310。复用器310是用于3个符号滤波器的图2的输入选择器283的示例。每个复用器310具有连接到寄存器282中的三个的三条输入线和连接到一个线程处理单元115的一条输出线。更具体地，每个复用器310可以经由其输入线从寄存器阵列282的相应子集将数据元素经由其输出线传递到其对应的线程处理单元115。例如，复用器M1可将来自寄存器阵列282的对应子集(包括R0、R1和R2)的数据传送到其对应的线程处理单元115，且复用器M2可将来自寄存器阵列282的对应子集(包括R1、R2和R3)的数据传送到其对应的线程处理单元115。因此，两个相邻复用器310可以重复使用相同的输入数据元素。例如，M1和M2可以重复使用存储在R1和R2中的数据元素。

请注意，复用器310是输入选择器的示例。复用器310可根据第一控制信号(例如，M1从R0接收第一数据元素)经由其输入线接收来自寄存器阵列282的对应子集的数据，也可以根据第二控制信号(例如，M1从R1接收第二数据元素)经由其它输入线接收来自耦合到另一复用器310的寄存器阵列282的另一子集的数据。可以生成额外的控制信号以控制进一步的数据传输。在一个实施例中，这些控制信号由控制器120产生。例如，控制器120可以在三个时刻发送第一、第二和第三控制信号，以分别选择每个复用器310的第一、第二和第三输入线，以将三个不同的数据元素传送到复用器的对应线程处理单元115。以此方式，可选择每一复用器310的所有三个输入线，且可将三个对应寄存器282中的数据元素发送到对应线程处理单元115中。

当将复用器310的集合作为整体考虑时，第一控制信号自寄存器282选择N个连续数据元素(例如d(0)、d(1)、...、d(N-1))的第一序列，并将它们发送到N个线程处理单元115，其中每个线程处理单元115接收第一序列中的数据元素之一。类似地，第二控制信号自寄存器282选择N个连续数据元素(例如，d(1)、d(2)、...、d(N))的第二序列，并且将它们发送到N个线程处理单元115，其中每个线程处理单元115接收第二序列中的数据元素之一。第三控制信号自寄存器282选择N个连续数据元素(例如，d(2)、d(3)、...、d(N+1))的第三序列，并且将它们发送到N个线程处理单元115，其中每个线程处理单元115接收第三序列中的数据元素之一。在该示例中，第一序列和第二序列被移位一个数据元素，并且类似地，第二序列和第三序列被移位一个数据元素。也就是说，在该示例中，每个序列是其先前序列的移位版本(通过移位一个数据元素)。在替代实施例中，每个数据序列可以从其紧邻的前导数据序列移位多于一个数据元素。

从该示例可以看出，MSE 380的使用显着减少了存储器访问量。在该示例中，(N+2)个数据元素d(0)、d(1)、...、d(N+1)从存储器(即，本地存储器116或系统存储器170)加载到MSE 380中一次，且(N-1)个数据元素d(2)、d(3)、...、d(N)的子集由线程处理单元115重复使用。如果每个线程处理单元310对其用于3个符号滤波操作的源操作数执行每通道存储器访问，则这些数据元素d(2)、d(3)、...、d(N)中的每一个将从存储器读取三次。相反，MSE380将(N+2)个数据元素的整个数据序列加载到其寄存器282中一次，而不重复读取相同的数据元素。MSE 380随后可以经由其复用器310高效地从其寄存器282向线程处理单元发送适当的源操作数。

在图3的实例中，MSE 380包括比用于在数据阵列的边界(即，开始和结束)保持两个额外数据元素的复用器310的数目多两个寄存器282。在替代实施例中，寄存器282的数量可以与复用器310的数量相同或比复用器310的数量多(不同于2的预定数量)。这些额外数目的寄存器282(如果有的话)影响可在滤波操作中使用的符号的数目、可从一个复用器310移位到下一个的数据元素的数目以及可由线程处理单元115执行的其它操作。

图4示出了根据一个实施例的MSE 480的另一示例。MSE 480包括耦合到输入选择器410的阵列的MSE寄存器282的阵列。在一个实施例中，每个输入选择器410包括至少一个复用器，例如图3所示的复用器310。在替代实施例中，每个输入选择器410包括多个复用器、开关或路由元件以选择来自寄存器282的相应子集的输入，并将这些输入传送到对应的线程处理单元115。

在图4的实施例中，指令解码器210解码指令并将关于已解码指令的信息传递到控制器120。根据该信息，控制器120产生一个或多个控制信号至输入选择器410，以控制它们对数据阵列的操作，包括但不限于：移位、混移，选择和传递。为了简化图示，图4仅示出作为控制信号的接收方的最左边的输入选择器410；但是应当理解，所有输入选择器410直接或间接地从控制器120接收控制信号。

在一个实施例中，每个输入选择器410经由输入线耦合到寄存器阵列282的相应子集(例如，输入选择器S1耦合到RA1，输入选择器S2耦合到RA2等)用于接收输入。在图4中，划分寄存器阵列282的虚线表示一些子集的边界。此外，每个输入选择器410还耦合到其相邻输入选择器410中的一个或两个。两个相邻输入选择器410之间的连接(称为“选择器间连接”)可以是单向的或双向的。因此，每个输入选择器410可以不仅经由输入线从其对应的子集接收数据输入，而且还经由选择器间连接从寄存器阵列282的其他子集接收数据输入。例如，输入选择器S1可以从RA1以及RA2、RA3、...等接收输入。

在一个实施例中，响应于传播控制信号，每个输入选择器410可以将数据元素从第一相邻输入选择器410传递到第二相邻输入选择器410(例如，数据可以从S3经由S2到S1，或从S4经由S2和S3到S1)。传播控制信号可以由控制器120或输入选择器410中指定的一个产生。

每个输入选择器410还经由输出线耦合到线程处理单元115的子集。在图4中，每个输出线可以由输入选择器410用于将输入数据传送到对应的线程处理单元115。输入数据可以从所选输入线或从所选择的选择器间连接接收。根据来自控制器120的控制信号，输入选择器410可以通过输入线或选择器间连接从寄存器中选择数据元素，并将该数据元素传送到线程处理单元115中的一些或全部，其中线程处理单元115连接到其输出。或者，输入选择器410可通过输入线和/或选择器间连接的组合从不同寄存器282中选择不同的数据元素，并将不同的数据元素传送到连接其输出的不同线程处理单元115。尽管图4中示出了特定数量的输入/输出线和选择器间连接，但是应当理解，在各种实施例中，每个输入选择器410可以连接到任何数量的输入/输出线和选择器间连接。

在一个实施例中，控制器120可以生成到输入选择器410的一个或多个控制信号，以控制它们对数据阵列的操作，包括但不限于：移位、混移、选择和传递。在一些实施例中，控制器120可以在由寄存器282的数量、输入线的数量和选择器间连接的数量所强加的限制内控制移位的方向和量，使得输出数据阵列是输入数据阵列的移位版本。控制器120还可以控制施加在输入数据阵列上的混移模式，以及选择哪些输入线和/或选择器间连接以使能到输出的数据传输。控制器120可以进一步命令输入数据阵列通过输入选择器410，使得输出数据阵列与输入数据阵列相同。

在图4中，以粗线示出了输入线和选择器间连接中的一部分，以指示这些线/连接被选择或启用。这些粗线示出了将输入数据阵列向左移位一个数据元素以产生输出数据阵列的示例。假定输入数据阵列是(d(0)、d(1)、d(2)、d(3)、...)，并且每个输入线将输入数据阵列的一个数据元素传送到输入选择器410。作为数据移位的结果，由输入选择器410生成的输出数据阵列是(d(1)、d(2)、d(3)、d(4)...)。如图3的示例中所示，K个符号滤波操作可以通过移位操作和通过(pass-through)操作的组合来实现。在数据移位的示例中，传播控制不被使能。

图5A和图5B根据一个实施例结合矩阵乘法示出了MSE 480的数据传输模式的示例。图5A示出了被乘数从所选输入线(示出为连接到S1的粗输入线)到所有输入选择器410的输出线的传播。例如，如果被乘数是(d(0)、d(1)、...、d(7))，输入选择器S1可以通过其所选择的输入线接收这些被乘数，并且相同的被乘数可以通过S1到达其输出线，并传播到其他输入选择器410以到达它们各自的输出线。在乘法的不同阶段，可以选择不同组的输入线。图5B示出了乘法器的传播：每个乘法器从输入选择器410的一个选定输入线(示为粗输入线)传送到该输入选择器410的所有输出线。在乘法的不同阶段，可以选择每个输入选择器410的不同的行。

除了上述滤波操作和矩阵乘法之外，MSE(例如，图1和图2的MSE 180，图3的MSE380和图4的MSE 480)可以具有其他信号、图像、多媒体处理应用，包括但不限于计算机视觉中的对象发现和模式识别。一种用于对象发现的已知技术将滑动窗应用于图像以检测给定对象(例如，人)的存在。滑动窗口可以具有固定尺寸，并且可以在图像的任何方向上滑动，诸如水平，垂直或对角线方向。滑动窗口的每个实例(即，图像的特定位置处的滑动窗口)可以被分配给线程处理单元115进行处理，并且滑动窗口的不同实例可以被分配给平行地不同的线程处理单元115。在一个实施例中，当滑动窗口在一个方向上滑动以创建滑动窗口实例的序列时，连续滑动窗口实例具有大量重叠的数据元素。因此，提供给第一线程处理单元的数据元素(源操作数)可以由第二线程处理单元重复使用。MSE可以用于向线程处理单元115提供数据元素，包括重叠的数据元素，而不具有由线程处理单元115的每通道访问引起的重复的相同数据访问。

例如，定向梯度的直方图(histogram of oriented gradients，简写为HOG)是在计算机视觉和图像处理中使用的用于对象发现的目的的特征描述符。该技术对图像的局部部分中的梯度定向的发生进行计数。图像被分成称为单元的小连接区域。对于每个单元内的像素，编译梯度方向的直方图。描述符是这些直方图的级联。HOG计算的第一步是梯度值的计算。一种用于计算梯度值的常用方法是在水平和垂直方向中的一个或两个方向上应用离散导数掩模(discrete derivative mask)。该掩模是滤波掩模，其被应用于图像单元以过滤图像的颜色或强度数据。因此，类似于对于K个符号滤波器的上述滤波操作，MSE也可以用于HOG计算的梯度值的计算。

如前所述，计算设备100响应主机150或计算设备100的指令集架构(instructionset architecture，简写为ISA)中定义的指令来执行操作。在一个实施例中，滤波器指令指定指示数据阵列的基本地址的第一操作数，指示数据阵列的大小的第二操作数，以及指示访问顺序(例如，线性、块、3D等)的第三操作数。以下伪代码提供了具有3个符号滤波器的滤波操作指令的示例：

//定义mse缓冲区大小

_local float mse[workgroup_size+2]；

*mse＝(image[f(workgroup_ID)],workgroup_size+2,linear)；

//定义如下：

//1.由于访问是基于工作组的，因此以workgroup_ID作为在存储器中的起始地址

//2.访问大小

//3.访问顺序(线性/块/3D/...)

//使用mse计算结果

float result＝(mse[woki_id]+mse[woki_id+1]+mse[woki_id+2])/3；

在一个实施例中，ISA还为可以利用上述MSE进行对象发现、矩阵乘法等的其他操作定义指令。

图6是示出根据一个实施例的由计算设备100执行的方法600的流程图。在一个实施例中，方法600可以由并行计算系统或设备(例如图1的计算设备100)执行。在图3和图4中提供了输入选择器的示例；也可以使用其他类型的输入选择器。

方法600开始于计算设备100接收识别数据元素的阵列作为输入的指令(步骤610)。计算设备100响应于指令将数据元素的阵列从存储器缓冲器加载到寄存器的阵列中(步骤620)。寄存器阵列耦合到输入选择器阵列，并且每个输入选择器经由输入线耦合到寄存器的相应子集。

根据第一控制信号，每个输入选择器将至少第一数据元素从寄存器的相应子集传送到一个或多个对应的线程处理单元(步骤630)。根据第二控制信号，每个输入选择器将来自寄存器的另一子集的至少第二数据元素传送到一个或多个对应的线程处理单元(步骤640)，所述另一子集经由其他输入线耦合到另一输入选择器。如上所述，不仅从输入选择器的相应子寄存器集合，而且从另一输入选择器的相应子寄存器集合选择输入的能力促进了数据重用，并减少了从存储器重复加载相同的数据的不必要的存储器流量。

在上面的描述中，假设N个线程处理单元被分配以并行执行滤波操作。在OpenCL中，分配给每个线程处理单元的计算任务称为工作项，并且共享数据和同步障碍的相关工作项形成工作组。通常，同一工作组的工作项使用相同的程序计数器(PC)并形成单批处理；因此，这些工作项将在锁步中逐步(step through)执行程序的指令。在某些情况下，工作组可能包含比计算单元中的线程处理单元数少的工作项。或者，工作组可以包含比计算单元中的线程处理单元数量更多的工作项，并且这些工作项的数量不能被线程处理单元的数量均分(evenly divisible)。在传统系统中的任何这些情况下，不能与工作项匹配的线程处理单元将被保留未使用。当并行计算中未使用某些线程处理单元时，将浪费系统资源。

根据本发明的一个实施例，计算单元中的线程处理单元被划分为多个批处理。分配到批处理的工作组数量(M)是可配置的。假设计算单元中的线程处理单元数为N，批处理中的线程处理单元数为P。M的取值范围为1≤M≤P。对于M个工作组中的每个工作组，工作组中可并行执行的工作项数为(P/M)，小于工作组的大小(即工作组中工作项的总数)。因此，它需要多次迭代来处理工作组中的所有工作项。

图7示出了根据一个实施例的具有可配置数量的工作组分配的计算单元710。作为计算设备700的一部分的计算单元710是图1的计算单元110的一个示例。根据图7，计算单元710包括附加的和/或替代的元件，其在图1的实施例中未示出。在该实施例中，计算单元710支持多个PC批处理和每个批处理中的可配置数量的工作组。计算单元710包括控制器单元750，在一个实施例中，控制器单元750可以执行图1的控制器120的操作。附加地或替代地，控制器单元750可以根据编译器701所决定的批处理设置720将M个工作组分配给批处理。在下面的描述中，跨不同批处理使用相同的可配置值M。在替代实施例中，不同批处理可以用不同数量的工作组分配。

在一个实施例中，可以在编译时决定可配置数量M(如图7所示)，或者在任务调度时决定可配置数量M。例如，编译器701可以分析指示要执行的任务的维度的输入NDRange，并且决定哪些工作组可以被打包到同一批处理中。可以使用的一个标准是具有相同同步障碍的工作组可以被打包到同一批处理中(即可配置批处理设置是从多个工作组的编译时的分析决定的)。如果M个工作组被打包到一个批处理中，则这些工作组可以根据相同的程序计数器760(图7中分别标示为CP0和PC1)执行指令。

为了说明的简单起见，图7在计算单元710中仅示出了一个计算单元710和两个批处理。应当理解，可以使用多于一个计算单元710和多于两个批处理来在一个输入中处理所有的工作组。每个批处理包括由批处理中的所有工作组、P个线程处理单元115、溢出存储器730和本地存储器116共享的程序计数器760。属于同一工作组的工作项可以使用本地存储器116用于共享数据。在一个实施例中，控制器单元750管理用于计算单元110中的不同批处理和批处理中的不同工作组的本地存储器116的分配和访问。

属于同一工作组的工作项也可以使用溢出存储器730来存储中间数据或上下文，例如当遇到同步障碍并且暂时中止(例如，通过进入等待状态)时。当该工作组中的所有工作项达到同步障碍时，等待状态结束。在一个实施例中，当批处理中的给定工作组遇到同步障碍时，给定工作组中的工作项可以被保存在溢出存储器730中，并且给定工作组中的其余工作项将被处理。假设分配给给定工作组的线程处理单元115的数量等于K(其中K＝P/M)，在一个实施例中，在给定的工作组可以移动通过同步障碍之前，K个线程处理单元115可循环遍历给定的工作组中所有剩余的工作项。在替代实施例中，批处理中的所有P个线程处理单元115可以专用于处理给定工作组中的所有其余工作项，以便整个给定工作组快速遇到(quickly bring the entire given workgroup to)同步障碍(synchronizationbarrier)。由于同步障碍而暂停的批处理中的工作项可以被暂时保存在溢出存储器730中。

图8A和图8B示出根据一个实施例的工作组分配到批处理中的示例。在这些示例中，四个工作组820(WG0、WG1、WG2、WG3)被分配给两个批处理(批处理0和批处理1)，其中每个工作组820包含多个工作项810。每个批处理包括两个工作组820，并且具有并行处理四个工作项的能力。图8A示出了并行处理的前四个工作项目，图8B示出了并行处理的下四个工作项目。在一些实施例中，不同工作组中的工作项可以是不同步的；因此，在一些实施例中，不同的工作组可以以不同的步速(pace)前进。然而，由于共享程序计数器，同一批处理中的工作组以相同的步速前进。

图9示出了根据一个实施例的由具有用于并行计算的可配置工作组分配的计算设备执行的方法900的流程图。在一个实施例中，方法900可以由并行计算系统或设备(诸如图7的计算设备700)执行。

方法900开始于控制器单元750将工作组分配给批处理的集合(步骤910)。至少一个批处理(“给定批处理”)被分配了共享程序计数器的M个工作组，M是根据可配置批处理设置决定的正整数。给定批处理包括一组线程处理单元。该组线程处理单元并行地执行M个工作组中的每个工作组中的工作项的子集(步骤920)。响应于M个工作组中的一个或多个遇到同步障碍的发现，M个工作组的中间数据被存储在溢出存储器中，并且M个工作组中的下一个工作项的子集被加载到线程处理单元中以并行执行(步骤930)。

前文已经描述了并行计算系统。并行计算系统可以克服常规硬件平台的限制，以实现高效的数据访问和重用，以及工作组的灵活分配。因此，可以提高系统的总体能量和计算效率。

已经参考图1、3、4和7的示例性实施例描述了图6和图9的流程图的操作。然而，应当理解，除了参考图1、3、4和7讨论的那些实施例实施图6和图9的流程图的操作之外，也可以参考图1、3、4和7讨论的实施例执行不同于这些流程图的操作。虽然图6和9的流程图示出了本发明的某些实施例执行的操作的特定顺序，但是应当理解，这种顺序是示例性的(例如，替代实施例可以以不同顺序执行操作，组合某些操作，重叠某些操作等)。

以上所述仅为本发明的较佳实施例，本领域相关的技术人员依据本发明的精神所做的等效变化与修改，都应当涵盖在权利要求内。

Claims (22)

1.一种执行多个并行计算的计算设备，其特征在于，所述计算设备包含：

多个线程处理单元；以及

存储器混移引擎，耦接于所述多个线程处理单元，所述存储器混移引擎包含：

多个寄存器组成的寄存器阵列，储存自存储器缓冲器获得的多个数据元素的阵列；以及

多个输入选择器组成的输入选择器阵列，每一输入选择器经由多个输入线耦接于所述多个寄存器的对应子集，并经由一个或多个输出线耦接于一个或多个对应线程处理单元，

其中所述存储器混移引擎；

响应于施加到所述多个输入选择器的第一控制信号，将所述多个寄存器的第一子集中的所述多个数据元素的第一序列并行发送到所述多个线程处理单元，其中每个线程处理单元接收相应的一个数据元素，以及

响应于施加到所述多个输入选择器的第二控制信号，将所述多个寄存器的第二子集中的所述多个数据元素的第二序列并行发送到所述多个线程处理单元，其中每个线程处理单元接收相应的一个数据元素，其中所述第二序列是所述第一个序列的移位版本。

2.根据权利要求1所述的执行多个并行计算的计算设备，其特征在于，每一输入选择器包含复用器，经由多个输入线耦接于所述多个寄存器的所述对应子集，并经由一个输出线耦接于一个线程处理单元。

3.根据权利要求1所述的执行多个并行计算的计算设备，其特征在于，每一输入选择器更经由一个或多个选择器间连接耦接于至少一个相邻输入选择器，并自所述至少一个相邻输入选择器接收一个或多个所述数据元素。

4.根据权利要求1所述的执行多个并行计算的计算设备，其特征在于，响应传播控制信号，每一输入选择器更自第一相邻输入选择器传送数据元素至第二相邻输入选择器。

5.根据权利要求1所述的执行多个并行计算的计算设备，其特征在于，每一输入选择器自多个不同寄存器选择多个不同数据元素，并将所述多个不同数据元素传送至多个不同线程处理单元。

6.根据权利要求1所述的执行多个并行计算的计算设备，其特征在于，每一输入选择器自所述多个寄存器选择一个数据元素，并将所述数据元素传送至多个不同线程处理单元。

7.根据权利要求1所述的执行多个并行计算的计算设备，其特征在于，更包含：

存储器混移控制器，利用一个或多个控制信号控制所述输入选择器阵列，以执行包含移位、混移、选择以及传递所述多个数据元素的阵列的至少一个操作。

8.一种执行多个并行计算的计算设备，其特征在于，所述计算设备包含：

控制单元，分配多个工作组至一组批处理；以及

所述组批处理，耦接于所述控制单元，每一批处理包含：

程序计数器，由分配给所述批处理的M个工作组共享，其中M是依据可配置批处理设置决定的正整数；

一组线程处理单元，并行执行所述M个工作组中的每一个工作组中的多个工作项的子集；以及

溢出存储器，当所述M个工作组中的一个或多个工作组遇到同步障碍时，储存所述M个工作组的中间数据。

9.根据权利要求8所述的执行多个并行计算的计算设备，其特征在于，所述可配置批处理设置是从所述多个工作组的编译时的分析决定的。

10.根据权利要求8所述的执行多个并行计算的计算设备，其特征在于，所述M个工作组中的每一个工作组包含比并行执行的所述多个工作项的所述子集多的工作项。

11.根据权利要求10所述的执行多个并行计算的计算设备，其特征在于，当所述多个工作项的所述子集完成或者暂停时，该组线程处理单元更并行执行所述M个工作组中的所述多个工作项的下一子集。

12.一种执行多个并行计算的计算方法，利用多个线程处理单元执行，其特征在于，所述计算方法包含：

接收将多个数据元素组成的数据元素阵列识别为输入的指令；

响应所述指令，自存储器缓冲器加载所述数据元素阵列至多个寄存器组成的寄存器阵列，其中所述寄存器阵列耦接于多个输入选择器组成的输入选择器阵列，且每一输入选择器经由多个输入线耦接于所述多个寄存器的对应子集；

响应于施加到所述多个输入选择器的第一控制信号，将所述多个寄存器的第一子集中的所述多个数据元素的第一序列并行发送到所述多个线程处理单元，其中每个线程处理单元接收相应的一个数据元素，以及

响应于施加到所述多个输入选择器的第二控制信号，将所述多个寄存器的第二子集中的所述多个数据元素的第二序列并行发送到所述多个线程处理单元，其中每个线程处理单元接收相应的一个数据元素，其中所述第二序列是所述第一个序列的移位版本。

13.根据权利要求12所述的执行多个并行计算的计算方法，其特征在于，所述每一输入选择器包含复用器，经由多个输入线耦接于所述多个寄存器的所述对应子集，并经由一个输出线耦接于一个线程处理单元。

14.根据权利要求12所述的执行多个并行计算的计算方法，其特征在于，更包含：

经由一个或多个选择器间连接，所述每一输入选择器自至少一个相邻输入选择器接收一个或多个所述数据元素。

15.根据权利要求12所述的执行多个并行计算的计算方法，其特征在于，更包含：

响应传播控制信号，所述每一输入选择器自第一相邻输入选择器传送数据元素至第二相邻输入选择器。

16.根据权利要求12所述的执行多个并行计算的计算方法，其特征在于，更包含：

所述每一输入选择器自多个不同寄存器选择多个不同数据元素；以及将所述多个不同数据元素传送至多个不同线程处理单元。

17.根据权利要求12所述的执行多个并行计算的计算方法，其特征在于，更包含：

所述每一输入选择器自所述多个寄存器选择一个数据元素；以及

将所述数据元素传送至多个不同线程处理单元。

18.根据权利要求12所述的执行多个并行计算的计算方法，其特征在于，更包含：

利用一个或多个控制信号控制所述输入选择器阵列，以使所述多个输入选择器对所述数据元素阵列执行包含移位、混移、选择以及传递操作的至少一个操作。

19.一种执行多个并行计算的计算方法，利用多个线程处理单元执行，其特征在于，所述计算方法包含：

分配多个工作组至多个批处理组成的一组批处理，其中所述多个批处理中的至少一个被分配给共享程序计数器的M个工作组，其中M是依据可配置批处理设置决定的正整数；

通过一组线程处理单元并行执行所述M个工作组中的每一个工作组中的多个工作项的子集；以及

响应所述M个工作组中的一个或多个工作组遇到同步障碍的发现，将所述M个工作组的中间数据储存至溢出存储器中，并将所述M个工作组中的多个工作项的下一子集加载至所述多个线程处理单元中用于并行执行。

20.根据权利要求19所述的执行多个并行计算的计算方法，其特征在于，更包含：

从所述多个工作组的编译时的分析决定所述可配置批处理设置。

21.根据权利要求19所述的执行多个并行计算的计算方法，其特征在于，所述M个工作组中的每一个包含比并行执行的所述多个工作项的所述子集多的工作项。

22.根据权利要求21所述的执行多个并行计算的计算方法，其特征在于，更包含：

当所述多个工作项的所述子集完成或者暂停时，所述组线程处理单元并行执行所述M个工作组中的所述多个工作项的所述下一子集。

Images