一种基于合并路径的SpMV实现方法、装置及存储介质

技术领域

本发明涉及高性能数值计算技术领域，尤其涉及一种基于合并路径的SpMV实现方法、装置及存储介质。

背景技术

稀疏线性代数的高性能算法在许多应用领域都很重要，包括计算科学、图分析和机器学习。稀疏矩阵向量乘(SpMV)作为高性能计算中常用的一项基础操作，在需要高性能计算任务中往往是性能瓶颈，因此加速SpMV操作具有十分重要的意义。

图像处理单元(GPU)具有吞吐量大、并行性高的特点，在科学计算领域是一个很有吸引力的选择。为了尽可能地提高在GPU上SpMV的运算速度，SpMV的算法设计需要遵从GPU的结构特点：1.线程束(warp)是GPU的基本调度单位，warp内的线程一次只能执行相同的指令，因此若条件分支过多会导致线程束分化严重，造成性能下降；2.当GPU访问内存的时候会一次性读取连续的内存空间，若目标地址分散在多段内存空间时，GPU需要读取多段连续的内存空间才能完成数据读取，因此若要提高GPU内存访问效率应尽可能地访问连续的地址空间；3.CUDA提供了多种warp级的线程通信原语，使用这些原语可以不通过共享内存实现极低开销的线程通信，利用这些warp级通信原语可以更高效地完成warp累加操作。

目前CSR存储格式是稀疏矩阵的主流存储格式，虽然对于新型稀疏矩阵存储格式的研究已有许多，但是格式转换的开销仍然不可忽略，所以研究基于CSR存储格式的SpMV算法依旧具有重要的意义。

当前基于CSR存储格式运行于GPU平台的SpMV算法研究，前人已做过许多工作。Duane Merrill和Michael Garland提出的Merge-based算法通过合并路径的方式来为每个线程公平地分配负载，这种负载分配方式可以有效应对稀疏矩阵连续长行和连续空行的情况，这使得Merge-based算法能够在具有不同稀疏度的稀疏矩阵下都有不错的性能，在多种SpMV算法中脱颖而出。

但是Merge-based算法仍有许多不足，主要如下：1.Merge-based算法为了确定每个线程负责的合并路径范围，需要进行合并路径搜索，这会造成一定的性能开销；2.Merge-based算法的数据访问方式是先将非零元素加载到共享内存，再从共享内存读取数据给线程，当遇到稀疏矩阵规模较大时，共享内存容量不足将会限制Merge-based算法的性能；3.由于Merge-based算法是以线程为单位分配合并路径负载，所以它访问共享内存的方式是非连续读取，这会造成内存访问效率低；4.由于合并路径的“写回”与“累加”是不规则的，那么处于同一个warp内的线程同时执行“写回”操作或同时执行“累加”操作的概率非常低，这就会造成严重的线程束分化问题，在稀疏矩阵平均行长较短的情况下这个问题会更加严重；5.由于存在矩阵行跨多个线程的情况，Merge-based算法需要对累加和进行修复，所需要修复的累加和个数等于线程数，当kernel函数调用过多的线程时，修复累加和将造成不小的性能开销。

发明内容

为至少一定程度上解决现有技术中存在的技术问题之一，本发明的目的在于提供一种基于合并路径的SpMV实现方法、装置及存储介质。

本发明所采用的技术方案是：

一种基于合并路径的SpMV实现方法，包括以下步骤：

对参数进行初始化；

确定每个GPU block在合并路径中的起始坐标和结束坐标；

确定每个线程束在合并路径中的起始坐标和结束坐标；

让每个线程束处理各自合并路径上的元素；

修复线程束之间的累加和；

其中，设SpMV操作为y＝Ax，其中A为稀疏矩阵，x和y为稠密向量；稀疏矩阵A具有num_rows行num_cols列，具有nnz个非零元素；所述SpMV实现方法在进行合并路径的搜索、稀疏矩阵非零元素的累加、非零元素累加和的写回都是以线程束为单位进行操作。

进一步地，所述稀疏矩阵A采用CSR存储格式来保存矩阵，包括三个数组values、column_indices和row_offsets；

合并路径中向下移动表示非零元素的累加，合并路径中向右移动表示将累加和写回到稠密向量y中，这种累加操作和写回操作称为合并路径元素(item)的处理。

进一步地，所述对参数进行初始化，包括：

确定合并路径的总长度nnz+num_rows；

确定每个线程束需要处理的合并路径长度num_items_per_warp；

确定GPU block的个数block_num和每个GPU block包含的线程数thread_num；

定义结构体数组pairs，结构体数组pairs的每个元素都包含两个变量：key和val。

进一步地，所述确定每个block在合并路径中的起始坐标和结束坐标，包括：

采用基于warp的合并路径搜索算法来确定每个GPU block在合并路径中的起始坐标和结束坐标；

所述基于warp的合并路径搜索算法的工作方式如下：

将线程束内的线程平均分散到合并路径的对角线上，判断是否有线程位于合并路径上；

若有线程位于合并路径上，将位于合并路径上的线程的坐标作为搜索结果，完成搜索；

若没有线程位于合并路径上，根据本轮搜索结果缩小对角线搜索范围，再将线程束内线程平均分散到搜小范围后的对角线上，判断是否有线程位于合并路径上，若有，完成搜索；若没有，继续缩小对角线搜索范围，直到有线程位于合并路径上，将位于合并路径上的线程的坐标作为搜索结果。

进一步地，所述确定每个线程束在合并路径中的起始坐标和结束坐标，包括：

采用基于warp的合并路径搜索算法来确定每个线程束在合并路径中的起始坐标和结束坐标。

进一步地，所述让每个线程束处理各自合并路径上的元素，包括：

在处理合并路径上的非零元素时，线程束内的线程将连续覆盖在合并路径上，线程束在合并路径上每次前进warp_size步，直到线程束遍历完其所负责的合并路径；其中warp_size为线程束大小；

由于稀疏矩阵每行具有的非零元素个数是不定的，因此线程束在合并路径上前进warp_size步时，线程束内的每个线程需要对合并路径进行搜索，以确定每个线程在合并路径中的坐标以及该坐标是否为矩阵行末。

进一步地，所述让每个线程束处理各自合并路径上的元素，包括：

在处理合并路径上的非零元素时，线程束每前进一次都需要对非零元素进行累加并写回累加和；

线程束对非零元素的累加操作包括：

线程束从GPU的全局内存中批量读取数据，使用线程束级操作原语__shfl_up_sync()进行循环右移累加，将线程束内非零元素的值累加到对应的矩阵行末；

线程束将矩阵行末的累加和批量写回的GPU的全局内存；完成后，线程束在合并路径上前进warp_size步并开始下一轮的累加与写回；其中warp_size为线程束大小。

进一步地，所述修复线程束之间的累加和，包括：

线程束在遍历完所有的合并路径后，保存跨线程束的累加和以及其对应的矩阵行号到数组pairs内；

当所有的线程束都遍历完各自的合并路径后，开始修复线程束间累加和：对数组pairs中处于同一矩阵行的累加和进行相加，并将各矩阵行结果写回到稠密向量y中。

本发明所采用的另一技术方案是：

一种基于合并路径的SpMV实现装置，包括：

至少一个处理器；

至少一个存储器，用于存储至少一个程序；

当所述至少一个程序被所述至少一个处理器执行，使得所述至少一个处理器实现上所述的方法。

本发明所采用的另一技术方案是：

一种计算机可读存储介质，其中存储有处理器可执行的程序，所述处理器可执行的程序在由处理器执行时用于执行如上所述方法。

本发明的有益效果是：本发明以线程束为单位来进行数据操作(包括访存、累加操作)，这样可以大大降低线程束分化程度，解决了Merge-based算法数据非连续读取导致访存效率低的问题。另外，以线程束为单位来分配合并路径负载，可以减少合并路径搜索总次数，同时每次合并路径搜索可以使用warp_size个线程，大大减少了合并路径搜索的总时间开销。

附图说明

为了更清楚地说明本发明实施例或者现有技术中的技术方案，下面对本发明实施例或者现有技术中的相关技术方案附图作以下介绍，应当理解的是，下面介绍中的附图仅仅为了方便清晰表述本发明的技术方案中的部分实施例，对于本领域的技术人员而言，在无需付出创造性劳动的前提下，还可以根据这些附图获取到其他附图。

图1是本发明实施例中基于warp的合并路径搜索算法的示意图；

图2是本发明实施例中warp_merge_search函数的伪代码的示意图；

图3是本发明实施例中确定每个block在合并路径中的起始坐标和结束坐标的伪代码的示意图；

图4是本发明实施例中确定每个warp在合并路径中的起始坐标和结束坐标的伪代码的示意图；

图5是本发明实施例中“warp的前进”的示意图；

图6是本发明实施例中warp_merge_forward函数的伪代码的示意图；

图7是本发明实施例中“warp的计算”示意图；

图8是本发明实施例中在多轮“warp的计算”中确保跨warp的矩阵行能够正确累加非零元素的示意图；

图9是本发明实施例中warp_merge_compute函数的伪代码的示意图；

图10是本发明实施例中warp处理合并路径上的元素的伪代码的示意图；

图11是本发明实施例中修复warp间累加和的伪代码的示意图；

图12是本发明实施例中一种在GPU上更高效的基于合并路径的SpMV实现方法的步骤流程图。

具体实施方式

下面详细描述本发明的实施例，所述实施例的示例在附图中示出，其中自始至终相同或类似的标号表示相同或类似的元件或具有相同或类似功能的元件。下面通过参考附图描述的实施例是示例性的，仅用于解释本发明，而不能理解为对本发明的限制。对于以下实施例中的步骤编号，其仅为了便于阐述说明而设置，对步骤之间的顺序不做任何限定，实施例中的各步骤的执行顺序均可根据本领域技术人员的理解来进行适应性调整。

在本发明的描述中，需要理解的是，涉及到方位描述，例如上、下、前、后、左、右等指示的方位或位置关系为基于附图所示的方位或位置关系，仅是为了便于描述本发明和简化描述，而不是指示或暗示所指的装置或元件必须具有特定的方位、以特定的方位构造和操作，因此不能理解为对本发明的限制。

在本发明的描述中，若干的含义是一个或者多个，多个的含义是两个以上，大于、小于、超过等理解为不包括本数，以上、以下、以内等理解为包括本数。如果有描述到第一、第二只是用于区分技术特征为目的，而不能理解为指示或暗示相对重要性或者隐含指明所指示的技术特征的数量或者隐含指明所指示的技术特征的先后关系。

此外，在本发明的描述中，除非另有说明，“多个”是指两个或两个以上。“和/或”，描述关联对象的关联关系，表示可以存在三种关系，例如，A和/或B，可以表示：单独存在A，同时存在A和B，单独存在B这三种情况。字符“/”一般表示前后关联对象是一种“或”的关系。

本发明的描述中，除非另有明确的限定，设置、安装、连接等词语应做广义理解，所属技术领域技术人员可以结合技术方案的具体内容合理确定上述词语在本发明中的具体含义。

术语解释：

GPU block：线程块，简称block。

warp：线程束。

为了解决现有技术merge-based算法存在的缺点，本发明提供一种面向GPU更高效的基于合并路径的SpMV算法。目前基于CSR存储格式的SpMV算法中，Merge-based算法通过以合并路径分配负载的方式，在具有不同稀疏度的稀疏矩阵下都达到不错的性能。但是Merge-based算法以线程为单位的思路并不符合GPU的底层硬件以warp为单位的特性，这会造成合并路径搜索开销大、非连续访存、累加操作多、线程束分化等问题。因此，本发明提出一种以warp为单位基于合并路径的SpMV算法。本算法将合并路径负载平均分配给每个warp，每次合并路径搜索将调用warp来并行搜索，warp内线程以连续而非分散的方式遍历合并路径，warp内线程同时从内存读取连续数据或写回数据，本算法使用CUDA提供的warp级原语进行warp内循环右移实现同一矩阵行非零元素的累加。这种以warp为单位的基于合并路径的思路，能够减少合并路径搜索次数、减少每次合并路径搜索的时间开销、减少需要修复的累加和个数、提高内存访问效率、无需占用大量的GPU共享内存、缓解线程束分化的问题、减少累加操作总次数，从而提高整体SpMV性能。

设稀疏矩阵A具有num_rows行num_cols列，一共具有nnz个非零元素。该稀疏矩阵A采用CSR存储格式来保存矩阵，包括三个数组values、column_indices和row_offsets：

1)values，数组大小为nnz，按照顺序逐行记录矩阵中非零元素的值；

2)column_indices，数组大小为nnz，按照顺序逐行记录矩阵中非零元素所在的列；

3)row_offsets，数组大小为num_rows+1，记录矩阵每行第一个元素在数组values和column_indices中的索引，其中row_offsets[num_rows]＝nnz。

SpMV操作可以表示为y＝Ax，其中y和x为稠密向量。设所使用GPU的线程束大小为warp_size，即每个线程束(warp)拥有warp_size个线程。

本申请提出的算法基于合并路径来分配负载，在进行合并路径搜索、稀疏矩阵非零元素的累加、非零元素累加和的写回都是以warp为单位进行操作。合并路径中向下移动表示非零元素的累加，合并路径中向右移动表示将累加和写回到y中，这种“累加”和“写回”操作称为合并路径元素(item)的处理。

本实施例提供一种面向GPU更高效的基于合并路径的SpMV实现方法，包括以下步骤：

S1、初始化。初始化过程需要确定每个warp需要处理的合并路径长度。

在本实施例中，初始化过程包括：

S11、计算合并路径总长度(nnz+num_rows)；

S12、确定每个warp需要处理的合并路径长度(num_items_per_warp)；

S13、确定kernel函数的线程规模，即GPU block的个数(block_num)和每个GPUblock包含的线程数(thread_num)，thread_num必须为warp_size的整数倍，则每个GPUblock拥有的线程束数量(warp_num)等于thread_num/warp_size；

S14、定义结构体数组pairs，pairs的大小为block_num*warp_num，数组pairs的每个元素都包含两个变量：key和val。为数组pairs分配全局内存空间，该数组将用于修复跨多个warp的矩阵行。

S2、确定每个block在合并路径中的起始坐标和结束坐标。

作为一种可选的实施方式，使用基于warp的合并路径搜索算法确定每个block和每个warp在合并路径中的起始坐标和结束坐标。所述基于warp的合并路径搜索算法需要调用一个warp来进行合并路径搜索，搜索过程为：首先将warp内的线程平均分散到合并路径的对角线上，然后判断是否有线程位于合并路径上，若有则完成搜索，若没有则根据本轮搜索结果缩小对角线搜索范围，再次将warp内线程平均分散到搜小范围后的对角线上，判断是否有线程位于合并路径上，如此循环直到有线程位于合并路径上，位于合并路径上的线程的坐标即为搜索结果。

首先说明基于warp的合并路径搜索算法(warp_merge_search函数)，其示意图如图1所示。首先根据item的个数确定确定对角线所处的位置。在第一轮循环时，属于同一个线程束的线程t

S3、确定每个warp在合并路径中的起始坐标和结束坐标。

确定每个warp在合并路径中的起始坐标和结束坐标。由于已经确定了block的起始坐标和结束坐标，那么在确定warp的起始坐标和结束坐标时就可以缩小范围合并路径的范围，并调用warp_merge_search函数来求解，其伪代码如图4所示。

S4、让每个warp处理各自合并路径上的元素。

每个warp处理各自合并路径上的元素。该步骤包括两种操作：“warp的前进”与“warp的计算”，分别对应warp_merge_forward函数和warp_merge_compute函数。

首先说明“warp的前进”，其示意图如图5所示。在确定了warp的起始坐标和结束坐标后，warp需要遍历它们之间的所有合并路径元素，且warp每次遍历会前进warp_size步，但由于矩阵每行用于的非零元素的个数是不定的，所以在warp前进后需要计算其坐标，以及确定该坐标是否处于矩阵行末(需要写回累加和)。warp_merge_forward函数的伪代码如图6所示。

接下来说明“warp的计算”。在调用warp_merge_forward函数后，可以确定warp内的每个线程各自对应的矩阵行号、非零元素的索引和是否属于矩阵行末，warp一次会处理warp_size个合并路径元素，包括“累加”与“写回”。“warp的计算”示意图如图7所示，图中阴影部分表示该线程处于矩阵行末。首先通过warp内右移实现同一行的累加，在图7的例子中warp_size等于16，为确保所有情况都能完成累加需要进行4次warp内右移累加，右移量分别为8、4、2、1，在右移累加的时候需要确保相加的两个线程处于同一矩阵行，若不处于同一矩阵行则不相加。完成右移累加后每个矩阵行的值都会累加到行末，此时再将该累加和写回到y向量中。当发生矩阵行跨warp的情况时，即warp的最后的一个线程所处的位置并不是矩阵行末时，需要将该行的累加和保存起来，在下一轮的“warp的计算”时首先加到warp的第一个线程处，以确保跨warp的矩阵行能够正确累加非零元素，示意图如图8所示。warp_merge_compute函数的伪代码如图9所示。

warp处理合并路径上的元素的过程就是“warp的前进”与“warp的计算”交替执行直到起始坐标和结束坐标之间的合并路径全部处理完，伪代码如图10所示。

S5、修复warp间累加和。

修复warp间累加和。当warp的结束坐标不是矩阵行末的时候，说明该矩阵行跨越了多个warp，因此需要将不同warp该矩阵行的累加和再进行相加，最后写回到y向量中。因此当warp的结束坐标不是矩阵行末的时候需要将最后的矩阵行的累加和暂存到pairs结构体数组中，pairs的结构体变量key和val分别用于记录矩阵行号与累加和，如图10中所示。当所有的warp都完成算法步骤S2、算法步骤S3和算法步骤S4之后就可以开始修复warp间累加和了，其伪代码如图11所示。

以下结合附图及具体实施例进行说明。

参见图12，本实施例还提供一种在GPU上更高效的基于合并路径的SpMV实现方法，设稀疏矩阵A具有num_rows行num_cols列，一共具有nnz个非零元素。该稀疏矩阵A采用CSR存储格式来保存矩阵，包括三个数组values、column_indices和row_offsets。SpMV操作可以表示为y＝Ax，其中y和x为稠密向量。设所使用GPU的线程束大小为warp_size。该方法具体包括以下步骤：

步骤1：初始化。确定kernel函数的线程规模，包括GPU block的个数以及每个GPUblock包含的线程数目(需为warp_size的整数倍)；确定每个warp需要处理的合并路径长度；定义包含两个结构体变量key和val的结构体数组pairs，并为其分配全局内存空间。

步骤2：使用基于warp的合并路径搜索算法(warp_merge_search函数)来确定每个block在合并路径中的起始坐标和结束坐标，如图2和图3所示。

步骤3：使用warp_merge_search函数来确定每个warp在合并路径中的起始坐标和结束坐标，如和图4所示。

步骤4：每个warp各自循环调用warp_merge_forward函数和warp_merge_compute函数，从而让warp遍历合并路径上的所有元素，并完成非零元素的累加与写回操作，最后将跨多个warp的矩阵行累加和暂存到pairs数组中，如图10所示。

步骤5：对保存在pairs数组中的累加和进行修复并将结果写回到y向量中，如图11所示。

综上所述，本申请的方法相对于现有技术，至少具有如下优点及有益效果：

(1)本申请以warp为单位来分配合并路径负载，这样可以减少合并路径搜索总次数，同时每次合并路径搜索可以使用warp_size个线程，大大减少了合并路径搜索的总时间开销。

(2)本申请减少了跨矩阵行累加和的修复次数，例如：若一共有128个线程且warp_size为32时，则Merge-based算法需要修复128个跨矩阵行累加和，而Warp-merge-based算法只需要修复4个跨矩阵行累加和。

(3)本申请的算法以warp为单位来进行数据操作(包括访存、累加操作)，这样可以大大降低线程束分化程度，并解决了Merge-based算法数据非连续读取导致访存效率低的问题。

(4)与Merge-based算法不同，本申请的算法使用warp内右移累加算法减少累加和所需的累加运算次数，例如：若warp_size为32，则Merge-based算法累加32个非零元素需要32次循环，本申请的算法只需要5次循环。

(5)相比于Merge-based算法，本申请的算法采用直接读非零元素的访存策略，无需占用大量共享内存空间来缓存非零元素。

本实施例还提供一种基于合并路径的SpMV实现装置，包括：

至少一个处理器；

至少一个存储器，用于存储至少一个程序；

当所述至少一个程序被所述至少一个处理器执行，使得所述至少一个处理器实现如图12所示方法。

本实施例的一种基于合并路径的SpMV实现装置，可执行本发明方法实施例所提供的一种基于合并路径的SpMV实现方法，可执行方法实施例的任意组合实施步骤，具备该方法相应的功能和有益效果。

本申请实施例还公开了一种计算机程序产品或计算机程序，该计算机程序产品或计算机程序包括计算机指令，该计算机指令存储在计算机可读存介质中。计算机设备的处理器可以从计算机可读存储介质读取该计算机指令，处理器执行该计算机指令，使得该计算机设备执行图12所示的方法。

本实施例还提供了一种存储介质，存储有可执行本发明方法实施例所提供的一种基于合并路径的SpMV实现方法的指令或程序，当运行该指令或程序时，可执行方法实施例的任意组合实施步骤，具备该方法相应的功能和有益效果。

在一些可选择的实施例中，在方框图中提到的功能/操作可以不按照操作示图提到的顺序发生。例如，取决于所涉及的功能/操作，连续示出的两个方框实际上可以被大体上同时地执行或所述方框有时能以相反顺序被执行。此外，在本发明的流程图中所呈现和描述的实施例以示例的方式被提供，目的在于提供对技术更全面的理解。所公开的方法不限于本文所呈现的操作和逻辑流程。可选择的实施例是可预期的，其中各种操作的顺序被改变以及其中被描述为较大操作的一部分的子操作被独立地执行。

此外，虽然在功能性模块的背景下描述了本发明，但应当理解的是，除非另有相反说明，所述的功能和/或特征中的一个或多个可以被集成在单个物理装置和/或软件模块中，或者一个或多个功能和/或特征可以在单独的物理装置或软件模块中被实现。还可以理解的是，有关每个模块的实际实现的详细讨论对于理解本发明是不必要的。更确切地说，考虑到在本文中公开的装置中各种功能模块的属性、功能和内部关系的情况下，在工程师的常规技术内将会了解该模块的实际实现。因此，本领域技术人员运用普通技术就能够在无需过度试验的情况下实现在权利要求书中所阐明的本发明。还可以理解的是，所公开的特定概念仅仅是说明性的，并不意在限制本发明的范围，本发明的范围由所附权利要求书及其等同方案的全部范围来决定。

所述功能如果以软件功能单元的形式实现并作为独立的产品销售或使用时，可以存储在一个计算机可读取存储介质中。基于这样的理解，本发明的技术方案本质上或者说对现有技术做出贡献的部分或者该技术方案的部分可以以软件产品的形式体现出来，该计算机软件产品存储在一个存储介质中，包括若干指令用以使得一台计算机设备(可以是个人计算机，服务器，或者网络设备等)执行本发明各个实施例所述方法的全部或部分步骤。而前述的存储介质包括：U盘、移动硬盘、只读存储器(ROM，Read-Only Memory)、随机存取存储器(RAM，Random Access Memory)、磁碟或者光盘等各种可以存储程序代码的介质。

在流程图中表示或在此以其他方式描述的逻辑和/或步骤，例如，可以被认为是用于实现逻辑功能的可执行指令的定序列表，可以具体实现在任何计算机可读介质中，以供指令执行系统、装置或设备(如基于计算机的系统、包括处理器的系统或其他可以从指令执行系统、装置或设备取指令并执行指令的系统)使用，或结合这些指令执行系统、装置或设备而使用。就本说明书而言，“计算机可读介质”可以是任何可以包含、存储、通信、传播或传输程序以供指令执行系统、装置或设备或结合这些指令执行系统、装置或设备而使用的装置。

计算机可读介质的更具体的示例(非穷尽性列表)包括以下：具有一个或多个布线的电连接部(电子装置)，便携式计算机盘盒(磁装置)，随机存取存储器(RAM)，只读存储器(ROM)，可擦除可编辑只读存储器(EPROM或闪速存储器)，光纤装置，以及便携式光盘只读存储器(CDROM)。另外，计算机可读介质甚至可以是可在其上打印所述程序的纸或其他合适的介质，因为可以例如通过对纸或其他介质进行光学扫描，接着进行编辑、解译或必要时以其他合适方式进行处理来以电子方式获得所述程序，然后将其存储在计算机存储器中。

应当理解，本发明的各部分可以用硬件、软件、固件或它们的组合来实现。在上述实施方式中，多个步骤或方法可以用存储在存储器中且由合适的指令执行系统执行的软件或固件来实现。例如，如果用硬件来实现，和在另一实施方式中一样，可用本领域公知的下列技术中的任一项或他们的组合来实现：具有用于对数据信号实现逻辑功能的逻辑门电路的离散逻辑电路，具有合适的组合逻辑门电路的专用集成电路，可编程门阵列(PGA)，现场可编程门阵列(FPGA)等。

在本说明书的上述描述中，参考术语“一个实施方式/实施例”、“另一实施方式/实施例”或“某些实施方式/实施例”等的描述意指结合实施方式或示例描述的具体特征、结构、材料或者特点包含于本发明的至少一个实施方式或示例中。在本说明书中，对上述术语的示意性表述不一定指的是相同的实施方式或示例。而且，描述的具体特征、结构、材料或者特点可以在任何的一个或多个实施方式或示例中以合适的方式结合。

尽管已经示出和描述了本发明的实施方式，本领域的普通技术人员可以理解：在不脱离本发明的原理和宗旨的情况下可以对这些实施方式进行多种变化、修改、替换和变型，本发明的范围由权利要求及其等同物限定。

以上是对本发明的较佳实施进行了具体说明，但本发明并不限于上述实施例，熟悉本领域的技术人员在不违背本发明精神的前提下还可做作出种种的等同变形或替换，这些等同的变形或替换均包含在本申请权利要求所限定的范围内。