CN101398753A - 用于执行扫描运算的系统、方法及计算机程序产品 - Google Patents

Abstract

本发明提供一种用于有效地执行扫描运算的系统、方法及计算机程序产品。在使用中，通过利用并行处理器架构来遍历元阵列。此并行处理器架构包含各自能够物理上并行执行预定数目的线程的多个处理器。出于效率目的，可执行所述处理器中的至少一者的所述预定数目的线程以执行涉及所述元的数目的扫描运算，所述元的数目是所述预定数目的线程的函数(例如，倍数等)。

Description

用于执行扫描运算的系统、方法及计算机程序产品

技术领域

本发明涉及并行处理器架构，且更明确地说涉及使用并行处理器架构执行计算算法。

背景技术

并行处理器架构通常用于执行不同计算算法的宽广阵列。通常使用所述架构来执行的算法的实例是扫描运算(例如，“all-prefix-sums”运算等)。一个此扫描运算定义于表1中。

                   表1

       $[I,a_0,(a_0\⊕a_1),\·\·\·,(a_0\⊕a_1\⊕\·\·\·a_{n-1}]),$

具体来说，假设阵列[a₀，a₁，…，a_n-1]及“I”是算子的单位元，那么返回表1的阵列。例如，如果算子

是加法算子，那么对阵列[3 1 7 0 4 1 6 3]执行扫描运算将返回[0 3 4 11 11 15 16 22]依此类推。尽管以上实例中阐述为加法算子，但此算子可以是以两个运算数运算的二元结合算子。

为对具有大量元的阵列有效地执行此扫描运算，可以“树”样方式遍历所述元。例如，可将所述元视为“树叶”，可在第一层级处处理“树叶”以产生并临时存储包含第一元的和的第二层级元等。此后，可以类似方式处理所述第二层级元，且直到已达到根。

为适应使用并行处理器架构的此处理，将每一阵列元指派给处理器的特定线程。通常，存在有限数目的处理器，而处理器各自又具有有限数目的线程(线程的量通常远小于阵列元的数目)。此外，由于线程共享从一个层级到下一层级的数据，因此在移动到下一层级之前必须完全完成每一前面层级的处理，等等。

而，此在每一层级的处理时需要同步。换句话说，扫描运算必须等到线程的指派且在移动到下一层级之前完成对特定层级处的每一阵列元的处理。例如，假设1024个元正由能够处理1元/时钟循环的32个线程来运算，那么以上算法在移动到下一层级的处理之前必须等待32个时钟循环。在使用中，前述同步可能造成闲置线程及额外的等待时间。

发明内容

附图说明

图1显示根据本发明的一个实施例用于利用并行处理器架构有效地执行扫描运算的系统。

图2显示根据本发明的另一实施例用于对阵列执行“异或”扫描运算的方法。

图3图解说明根据本发明的另一实施例的实例性“异或”扫描运算。

图4显示根据本发明的再一实施例用于对较大的元阵列执行扫描运算的框架。

图5图解说明其中可实施各种先前实施例的各种架构及/或功能性的实例性系统。

具体实施方式

图1显示根据本发明的一个实施例用于利用并行处理器架构101有效地执行扫描运算的系统100。在本说明的上下文中，并行处理器架构可包含任一架构，只要其包含并行运算的两个或两个以上处理器102A-N。在一个实施例中，此并行处理器架构可采取以下形式：图形处理器[例如，图形处理单元(GPU)等]或中央处理器单元(CPU)或配备有图形处理能力的芯片组。然而，当然，本发明预期其中并行处理器架构还采取其它形式(例如，通用计算处理器等)的其它实施例。

如进一步所示，并行处理器架构的多个处理器各自能够物理上并行地执行预定数目的线程104。在一个实施例中，此对线程的物理执行是指能够同时物理上执行的线程的数目，此不同于以逻辑方式的执行(例如，使用时间分片技术等)。

作为选项，每一处理器的线程可以单指令多数据(SIMD)方式运算。换句话说，处理器的所有线程可同时对不同数据执行同一指令。在一个实施例中，可将以此方式运算的此一组线程称为“卷曲”。此外，预定数目的线程可称为对应处理器的“卷曲大小”。

在使用中，通过利用并行处理器架构来遍历元阵列。在本说明的上下文中，此阵列的元可包含能够经受扫描运算的任一组的值。例如，在一个实施例中，通常可由表达式[A，B，C…N]来表示所述阵列的值，其中所显示的值为数字值。当然，阐述此元阵列仅出于说明目的而无论如何不应被视为以任何方式加以限制。

在遍历阵列元期间，可执行扫描运算。在本说明的上下文中，扫描运算可指涉及所述阵列的当前元及至少一个先前元的任一运算(如果可用的话)。在一个实施例中，扫描运算可包含all-prefix-sums运算。将在描述图2中所图解说明的不同实施例期间阐述关于实例性all-prefix-sums运算的更多信息。当然，本发明预期其它扫描运算(例如，涉及更多或更少元及其它算子等)，只要所述扫描运算满足以上定义。

出于效率目的，可执行至少一个处理器的预定数目的线程来执行涉及所述元的数目的扫描运算(例如，前述“卷曲大小”等)，所述元的数目是预定数目的线程的函数。例如，在一个实施例中，可执行预定数目的线程以执行涉及所述元的数目是所述预定数目的倍数的扫描运算。在本说明的上下文中，线程的预定数目的前述倍数可包含任一整数(例如，1、2、3、4、5…N等)。在图1中所示实施例中，扫描运算可涉及等于线程的预定数目的元的数目(即，所述倍数等于1)。

在任何情况下，可给特定处理器的每一线程指派一元以用于执行相关扫描运算。为此，如果不能一起避免的话，那么也可减小与线程之间的同步相关联的处理。换句话说，作为以上设计的结果，可准确地给每一线程指派一个元以执行扫描运算，从而使得特定处理器的所有线程可同时终止。作为选项，可利用任选XOR运算或类似运算来遍历所述阵列的元，以提供额外的效率。

现在将关于各种任选架构及特征阐述更多说明性信息，按照用户的需要可以或可不利用所述各种任选架构及特征来实施前述框架。具体来说，将阐述至少一个额外实施例，其结合扫描运算使用“异或”来遍历元阵列。应极其注意，阐述以下信息仅出于说明目的，而不应被视为以任何方式加以限制。任一以下的特征可视需要并入，而并不排除所述的其它特征。

图2显示根据本发明的另一实施例用于对阵列P执行“异或”扫描运算的方法200。作为选项，本方法可在图1的功能性及架构的上下文中实施。例如，可通过并行处理器架构(例如图1中所示的并行处理器架构)中的特定处理器的线程跨越多个阵列元来实施本发明。然而，当然，可在所需的环境中实施所述方法(例如，在没有图1的设计准则的情况下等)。还应注意可在本说明期间应用前述定义。

如所示，所述方法通过将变量D设定为“1”初始化变量D而开始。参见运算202。当然，此初始化是任选的且如果存在的话，那么可以任一所需的方式来执行。接下来，所述方法在当回路203中继续直到变量D达到卷曲大小。参见决定204。

同样，此卷曲大小是指能够物理上并行运行于并行处理器架构的特定处理器上的预定数目的线程。此外，在当回路内同步可能未必是需要的。具体来说，通过将阵列元的数目限制为小于或等于卷曲大小，那么在线程中间未必需要同步。如早期所提及，此同步涉及其中扫描运算必须等待线程的指派的情况且在移动到下一层级之前完成对每一阵列元的处理等。为此，本锁定步骤设计可避免其中第一线程尚未完成写入到存储器中后续线程需要读取或写入的共享部分的情况等。

如不久将变得明了，变量D针对当回路203的每一迭代增加一倍。通过以此方式将变量D增加到两倍，便可将所述阵列处理为二元树。在此上下文中，变量D与此树的层级相关。

在使用当期间，变量D保持小于卷曲大小，条件分支继续进行，如运算206-208中所示。具体来说，首先确定涉及变量D及线程-局部变量idx的逐位“与”运算是否大于“0”。参见决定206。此线程-局部变量idx是指多个活动线程中间特定线程的全局索引。在一个实施例中，idx可包含在使用期间被指派给线程的局部变量。此线程-局部变量可通过硬件来指派且可进一步在寄存器中追踪/存储。

表2 图解说明针对D及idx的不同值的决定206的结果。

                   表2

       D＝1:(idx AND D)>0，when idx＝1，3，5，7etc.

       D＝2:(idx AND D)>0，when idx＝2，3，6，7etc.

       D＝4:(idx AND D)>0，when idx＝4，5，6，7，12，13，14，15etc.

       D＝8:(idx AND D)>0，when idx＝8，9，10，11，12，13，14，15，24，25etc.

如果逐位“与”运算每一决定206大于“0”，那么更新阵列P的特定元。具体来说，仅通过所述树的最低层级处的对应线程(如以上表2中所阐述)来更新奇数元。

在每一决定206逐位“与”运算大于“0”时，基于以下表达式#1更新阵列P的特定元P[idx]。

           表达式#1

     P[idx]+＝P[(idx OR(D-1))XOR D]

此阵列元P[idx]的值显示为涉及变量idx及(D-1)的值的逐位“或”运算以及此结果与变量D的值的逐位“异或”两者的函数。

表3 图解说明针对idx及D的各种值求和为P[idx]的各种元的概要。

                        表3

               针对以下值被求和为P[idx]的元：

       idx＝     D＝1      D＝2      D＝4    D＝8     D＝16

       0         -         -         -       -        -

       10        -         -         -       -        -

       2         -         1         -       -        -

       3         2         1         -       -        -

       4         -         -         3       -        -

       5         4         -         3       -        -

       6         -         5         3       -        -

       7        6          5         3       -        -

       8        -          -         -       7        -

       9        8          -         -       7        -

       10       -          9         -       7        -

       11       10         9         -       7        -

       12       -          -         11      7        -

       13      12          -         11      7        -

       14      -          13         11      7        -

       15      14         13         11      7        -

       16      -          -          -       -        15

              针对以下值被求和为P[idx]的元：

      idx＝    D＝1    D＝2    D＝4    D＝8    D＝16

      17       16      -       -       -       15

      18       -       17      -       -       15

      19       18      17      -       -       15

      20       -       -      19       -       15

      21       20      -      19       -       15

      22       -      21      19       -       15

      23       22     21      19       -       15

      24       -       -       -       23      15

      25       24      -      -        23      15

      26       -      25      -        23      15

      27       26     25      -        23      15

      28       -      -       27       23      15

      29       28     -       27       23      15

      30       -      29      27       23      15

      31       30     29      27       23      15

      etc.

将在涉及图3中所示的8-元阵列的不同实施例上下文中阐述另一运算实例的图解。

在运算208之后，变量D增加一倍。参见运算210。此后，运算在当回路中继续直到变量D不再小于卷曲大小。同样参见决定204。在一个实施例中，当回路的结束可导致本方法的终止。在此实施例中，结果可采取相容“异或”扫描的形式。

在另一实施例中，所述方法可视需要以运算212继续进行，其中如下所述实施表达式#2。

                          表达式#2

                    P[idx]＝P[idx]-oval，

                       where oval＝P[idx]

应注意表达式“卵形＝P[idx]”是在回路开始之前在运算202中执行。否则，P[idx]将具有在回路中计算的新值，从而导致不正确的结果(例如，全部为零)。

在使用中，表达式#2的计算可用于将相容“异或”(inclusive XOR)结果转换为相斥“异或”(exclusive XOR)结果。相斥扫描可指其中所述结果的每一元j是输入阵列中所有元(但不包含元j)加起来的和的扫描。另一方面，在相容扫描中，对所有元(包含元j)求和。如运算212中所述，可通过将所得阵列向右移位一个元并插入单位元而从相容扫描来产生相斥扫描。应注意相斥扫描可指其中所述结果的每一元j是输入阵列中所有元(但不包含j)加起来的和的扫描。另一方面，相容扫描是其中对所有元(包含j)求和的扫描。

在使用中，前述方法可在多个线程上并行执行，且每一卷曲内的所有线程计算等于卷曲大小的数目的元的扫描。使用逐位“异或”运算，所述方法通过遍历呈树形式的阵列来建造扫描运算的结果。在所述树的每一层级D处，所述方法计算2D与每一线程索引的较低D位的“异或”，以计算所述线程读取的地址。实践中，由于卷曲大小对于给定机器是固定的，因此以上方法中的当回路被展开。

表4中阐述可用于实施前述方法的实例性伪代码。当然，阐述此伪代码仅出于说明目的且无论如何不应被视为以任何方式加以限制。

表4

warpscan(array P)

{

  Thread-local variables：idx，oval，

  idx＝this thread’s global index amongall active threads

  oval＝P[idx]

  D＝1

  while(D<warp_size)

  if((idx AND D)>0)then

  P[idx]+＝P[(idx OR(D-1))XOR D]

  endif

  D＝D*2

  endwhile

  if(this is an exclusive scan)

  P[idx]-＝oval

}

在一个实施例中，可利用任何所需的编程框架来实施本方法。在一个实施例中，可将驱动器用于利用图形处理器来提供一般计算能力来实施此技术。可结合英伟达(NVIDIA)公司提供的CUDA^TM框架来提供此驱动器的实例。表5图解说明用于支持此实施方案的实例性代码。同样，应极其注意阐述此实施方案仅出于说明目的，而无论如何不应被视为以任何方式加以限制。

表5

__device__warpscan(float* ptr，bool isExclusive)

{

  int idx＝threadIdx.x；

  float oval＝ptr[idx]；

  if(idx&1)    ptr[idx]+＝ptr[(idx|  (1-1))    ^1]；

  if(idx&2)    ptr[idx]+＝ptr[(idx|  (2-1))    ^2]；

  if(idx&4)    ptr[idx]+＝ptr[(idx|  (4-1))    ^4]；

  if(idx&8)    ptr[idx]+＝ptr[(idx|  (8-1))    ^8]；

  if(idx&16)   ptr[idx]+＝ptr[(idx|(16-1))^16]；

  if(isExclusive)

  ptr[idx]-＝oval；

}

图3显示根据本发明的另一实施例的实例性“异或”扫描运算300。作为选项，本扫描可表示图2的方法的实例性运算。例如，图解说明涉及八个元的元阵列302的多个通过。

具体来说，第一通过304显示为涉及元1、3、5、7等的更新。所述元的选择可由图2的决定206来规定。参见(例如)上文表2，在D＝1时。此外，此替换的值是使用多个“异或”运算305求出，且可根据图2的运算208来计算。注意，例如，表达式#1。

运算以所示方式继续第二通过306及第三通过308。如进一步图解说明，第三通过的最终元310包含所述元阵列的所有元的和。

如图3中所图解说明，出于效率目的，“异或”运算提供在单个方向(例如，向下等)上的阵列的遍历。在本说明的上下文中，此单个方向遍历可指避免在相反反向上的遍历(此将需要额外的处理)的任何遍历。在图3的实施例的特定上下文中，此相反方向上的遍历将涉及阵列的向上搜索。当然，在一些实施例中，预期使用多方向遍历。

应注意阐述“异或”运算的使用仅出于说明目的。预期将其它算子(例如，减算子等)用于提供类似于表3中所阐述的功能性的其它实施例。在一些实施例中，可使用任何适合的遍历方案。

图4显示根据本发明的再一实施例的用于对较大元阵列执行扫描运算的框架400。作为选项，本框架可使用图1-3的功能性及特征来实施。然而，本框架当然可在任何所需的环境中实施。同样，可在本说明期间应用前述定义。

在本实施例中，可提供对于以先前图中所阐述的方式进行处理太大的元阵列402。具体来说，特定处理器的线程不足以适应所述阵列的元数目。在此情况下，可将所述阵列的值划分成多个块404。例如，此块的大小可包含等于能够通过特定处理器物理上并行运行(例如，卷曲等)的线程数目的数目的元。

为此，可给每一块指派卷曲。通过此设计，可给所述卷曲的每一线程分配对应块的特定元。此外，多个处理器可各自处理相关联块的元以用于执行扫描运算。例如，参见图2。

可接着将每一块的扫描的结果存储于辅助阵列406中，以供在完成扫描运算中使用。具体来说，可将每一块的最后元存储于此辅助阵列中。进一步来说，可接着扫描此辅助阵列的元以用于产生扫描结果的额外阵列408。可接着将此扫描结果加到原始经扫描块404。明确来说，在项406/408表示相容扫描的情况下，可以所示的方式将扫描结果i加到原始经扫描块i+1的每一元。为此，提供经扫描元的最终阵列。尽管未显示，但在项406/408表示相斥扫描的情况下，可将扫描结果i加到原始经扫描块i的每一元。当然，尽管本实例涉及加法算子，但在各种实施例中此算子可包含但未必限定于乘法、最大、最小、逐位“与”/“或”等。

因此，单卷曲“异或”扫描运算可扩展到较大阵列。总之，可通过将扫描子阵列Ai的结果的最后元加到扫描子阵列Aj的结果的每一元，从许多非重叠阵列部分(例如，A的子阵列的前置和)来计算大阵列A的前置和。因此可利用此特性设计可使用B个线程扫描B个元的算法，其中B是卷曲大小的倍数。

在一个特定使用实例中，每一线程i可从装置存储器加载一个元且将其存储于共享存储器中的阵列P的位置i中。接着，可通过所有线程在阵列P上运行以上算法。此产生现在含有warp_size元的B/warp_size子阵列的P，每一子阵列含有所述输入的对应元的前置和。通过其对应卷曲w将这些子阵列扫描的每一者的最后元拷贝到另一共享阵列Q(仅具有B/warp_size)元的元w。接着扫描此阵列。最后，来自卷曲w＝取最小(i/warp_size)的每一线程i将阵列Q的元w加到阵列P的元i。阵列P因此含有输入阵列的完整前置扫描。

同样，由于块大小经设定而包含等于卷曲大小的数目的元(即，能够通过特定处理器物理上并行运行的线程的数目)，在块的扫描内未必需要同步。然而，尽管鉴于此设计可减小同步，但可在各种点利用某一同步。例如，可在对阵列的不同部分(例如，块等)执行扫描运算的线程中间执行同步。

表6阐述可用于实施图4的前述框架的实例性伪代码。如所示，在各种势垒点处提供同步。当然，阐述此伪代码仅出于说明目的且无论如何不应被视为以任何方式加以限制。

表6

scan(array P)

{

  Thread-local variables：idx，ival，oval

  idx＝this thread’s global index among all active threads

  ival＝P[idx]

  BARRIER

  warpscan(P)

  oval＝P[idx]

  BARRIER

  if(idx is the last thread in a warp)then

  P[

idx/warp_size

]＝oval+ival

  endif

  BARRIER

  if(idx<warp_size)then

  p[idx]＝warpscan(p)

  endif

  BARRIER

  oval+＝P[

idx/warp_size

]

  P[idx]＝oval

}

应注意在一个实施例中，前述“卷曲扫描”函数可通过许多卷曲在表6的伪代码中同时运行，而非仅通过一个卷曲。

在以上伪代码中，术语“势垒”是指在任一线程可在往前继续进行之前所有线程应达到的势垒同步点。在各种实施例中，此可用于避免读取后写入(WAR)及写入后读取(RAW)数据危险。

类似于先前实施例，本技术可利用任何所需的编程框架来实施。在一个可能的实施例中，可通过驱动器结合前述CUDA^TM框架来提供前述功能性。表7图解说明用于支持此实施方案的实例性代码。

表7

__global__void scan(float*g_odata，float*g_idata)

{

  extern__shared__float p[]；

  int idx＝threadIdx.x；

  float ival＝g_idata[idx]；

  p[idx]＝ival；

  __syncthreads()；

  float oval＝warpscan(p)；

  __syncthreads()；

  if((idx & 31)＝＝31)p[idx>>5]＝oval+ival；

  __syncthreads()；

  if(idx<32)p[idx]＝warpscan(p)；

  __syncthreads()；

  oval+＝p[idx>>5]；

  g_odata[idx]＝oval；

    }

应注意前述扫描运算可并行用于各种应用，包含但不先于分类(例如，基数分类等)、词汇分析、串比较、多项式求值、流紧缩、构造直方图及数据结构(例如，图形、树、经求和的区域表等)。当然，所述应用阐述为实例，而且预期其它实例。

图5图解说明可实施各种先前实施例的各种架构及/或功能性的实例性系统500。如所示，提供包含连接到通信总线502的至少一个主机处理器501的系统500。所述系统也包含主存储器504。控制逻辑(软件)及数据存储于可采取随机存取存储器(RAM)形式的主存储器中。

所述系统也包含图形处理器506及显示器508，即计算机监视器。在一个实施例中，所述图形处理器可包含多个着色器模块、光栅化模块等。每一前述模块甚至可设置于单个半导体平台上以形成图形处理单元(GPU)。

在本说明中，单个半导体平台可指单独整体式基于半导体的集成电路或芯片。应注意，术语“单个半导体平台”还可指具有增强连接性的多芯片模块，其模拟芯片上运算且对使用常规中央处理单元(CPU)及总线实施方案做出显著改进。当然，也可按照用户的需要，单独地或者以半导体平台的各种组合形式设置各种模块。

所述系统还可包含辅助存储装置510。所述辅助存储装置包含(例如)硬盘驱动机及/或可拆卸存储装置驱动机(其表示软盘驱动机、磁带驱动机、光盘驱动机等)。可拆卸存储装置驱动机以众所周知的方式从可拆卸存储单元读取及/或向可拆卸存储单元写入。

计算机程序或计算机控制逻辑算法可存储在主存储器及/或辅助存储装置中。在执行时，所述计算机程序启用系统以执行各种功能。存储器、存储装置及/或任何其它存储装置均是计算机可读媒体的可能实例。

在一个实施例中，各种先前图的架构及/或功能性可在以下上下文中实施：主机处理器、图形处理器、能够具有主机处理器及图形处理器两者的能力的至少一部分的集成电路(未显示)、芯片组(即，经设计以作为用于执行相关功能的单元而工作及出售的集成电路群组等)及/或用于所述事情的任何其它集成电路。此外，在一个可能的实施例中，各种先前图的所述与扫描有关的功能性可在驱动器512的控制下在任一前述集成电路中加以实施。

进一步来说，各种先前图的架构及/或功能性可在以下上下文中实施：通用计算机系统、电路板系统、专用于娱乐目的的游戏控制台系统、专用系统及/或任何其它所需的系统。例如，所述系统可采取以下形式：桌上型计算机、膝上型计算机及/或任何其它类型的逻辑。进一步来说，所述系统可采取各种其它装置的心事，包含但不限于个人数字助理(PDA)装置，移动电话装置、电视机等。

此外，尽管未显示，但所述系统可出于通信目的而耦合到网络[例如，电信网络、局域网络(LAN)、无线网络、例如因特网的广域网络(WAN)、对等网络、电缆网络等)。

尽管上文已描述了各种实施例，但应了解，所述实施例仅以实例的方式而非限制的方式呈现。因此，优选实施例的广度及范围不应受限于任何上述实例性实施例，而应仅根据以下权利要求书及其等效内容来界定。

Claims (20)

1、一种方法，其包括：

通过利用包含多个处理器的并行处理器架构来遍历元阵列，所述多个处理器各自能够物理上并行地执行预定数目的线程；及

执行所述处理器中的至少一者的所述预定数目的线程，以执行涉及所述元的数目的扫描运算，所述元的数目是所述预定数目的线程的函数。

2、如权利要求1所述的方法，其中所述线程各自对不同数据执行单个指令。

3、如权利要求1所述的方法，其中所述并行处理器架构包含图形处理器。

4、如权利要求1所述的方法，其中所述扫描运算包含all-prefix-sums运算。

5、如权利要求1所述的方法，其中在单个方向上遍历所述元阵列。

6、如权利要求1所述的方法，其中利用“异或”运算来遍历所述元阵列。

7、如权利要求1所述的方法，其中所述函数包含倍数。

8、如权利要求7所述的方法，其中所述倍数为1。

9、如权利要求7所述的方法，其中所述倍数至少为2。

10、如权利要求9所述的方法，其中对所述阵列的多个部分执行所述扫描运算，所述多个部分各自包含等于所述预定数目的元数目。

11、如权利要求10所述的方法，其中所述阵列的所述部分不重叠。

12、如权利要求10所述的方法，其中在对所述部分中的第一者执行所述扫描运算的所述线程与对所述部分中的第二者执行所述扫描运算的所述线程中间执行同步。

13、如权利要求10所述的方法，其中存储对所述阵列的所述部分执行的所述扫描运算的结果。

14、如权利要求13所述的方法，其中使用对所述阵列的所述部分执行的所述扫描运算的所述结果来完成所述扫描运算。

15、一种体现于计算机可读媒体上的计算机程序产品，其包括：

用于通过利用包含多个处理器的并行处理器架构来遍历元阵列的计算机代码，所述多个处理器各自能够物理上并行执行预定数目的线程；及

用于执行所述处理器中的至少一者的所述预定数目的线程以执行涉及所述元的数目的扫描运算的计算机代码，所述元的数目是所述预定数目的线程的函数。

16、如权利要求15所述的计算机程序产品，其中所述计算机代码是能够利用图形处理器提供一般计算能力的驱动器的组件。

17、如权利要求15所述的计算机程序产品，且其进一步包括用于在单个方向上遍历所述元阵列的计算机代码。

18、如权利要求15所述的计算机程序产品，且其进一步包括用于利用“异或”运算来遍历所述元阵列的计算机代码。

19、一种系统，其包括：

并行处理器架构，其包含各自能够物理上并行执行预定数目的线程的多个处理器；及

驱动器，其与所述并行处理器架构连通以用于执行所述处理器中的至少一者的所述预定数目的线程以执行涉及阵列元的数目的扫描运算，所述阵列元的数目是所述预定数目的线程的函数。

20、如权利要求19所述的系统，其中所述并行处理器架构经由总线耦合到存储器。

Images