文摘gydF4y2Ba

Aho-Corasick (AC)算法是一种常用的多模式字符串匹配算法在计算机和网络安全与生物信息学中,和其他很多。为了满足高度要求计算需求强加给这些应用程序,实现高性能AC算法是至关重要的。在本文中,我们提出了一个高性能的并行化交流等许多核心加速器芯片从Nvidia图形处理单元(GPU), Intel Xeonφ。我们的并行化方法显著提高了缓存位置的AC分区组给定的字符串模式分成多个小套模式空间。使用多个模式集,密集模式匹配操作同时进行对整个输入文本数据。相比以前的方法输入数据的分区中多个线程而不是分区模式集,我们的方法极大地提高了性能。实验结果表明,我们的方法导致了Nvidia甘蓝型GPU加速2.73倍和2.00倍加速在Intel Xeonφ与前面的方法。我们的并行实现提供高达693 Gbps甘蓝型吞吐量性能。gydF4y2Ba

1。介绍gydF4y2Ba

最近,许多核心加速器芯片等图形处理单元(gpu)英伟达和AMD和英特尔的许多集成的核心(MIC)体系结构,其中,正变得越来越受欢迎。这些芯片是快速增长的影响在高性能计算(HPC)服务器市场,在前500名单,特别是。他们有大量的核心和多个线程每核心级别的缓存层次结构,大量的(> 5 GB)车载记忆,和> 1 Tflops峰值双精度运算每个芯片的性能。他们大多是利用协处理器和并行程序执行内核由主机CPU对提供的输入数据从主机内存到车载设备内存。使用许多核心的加速器,许多创新的性能改进已经报道了HPC应用程序和更多的仍然是。gydF4y2Ba

字符串匹配算法是一个重要的在计算机和网络安全gydF4y2Ba1gydF4y2Ba- - - - - -gydF4y2Ba6gydF4y2Ba,生物信息学gydF4y2Ba7gydF4y2Ba,gydF4y2Ba8gydF4y2Ba),等等。在许多字符串匹配算法,Aho-Corasick (AC) [gydF4y2Ba9gydF4y2Ba)是一种多模式字符串匹配算法,可以同时匹配模式对于一个给定的有限集合的字符串(或字典)。一个确定性有限自动机(DFA)首先对给定的一组模式字符串构造。那么模式匹配操作时对输入文本进行数据引用DFA。按顺序输入数据访问;因此,访问模式是完全可预测的。然而,DFA的访问是不规则的跳跃的有很多从一个状态到另一状态的DFA处理顺序输入字符时可能的匹配。模式字符串数量的增加,例如,几十的成千上万的病毒在电脑病毒扫描模式字符串(gydF4y2Ba1gydF4y2Ba),相应地在DFA数量的增加。大量的国家在DFA与不规则数据结构访问会导致穷人数据局部性和高缓存错过。因此,为了加快模式匹配和满足性能需求强加于交流,优化缓存位置是至关重要的。gydF4y2Ba

在本文中,我们开发一个高性能并行的AC字符串匹配算法,显著提高了缓存位置上不定期的DFA访问许多核心加速器芯片如Nvidia Tesla甘蓝型GPU和Intel Xeonφ。先前的研究并行化交流(gydF4y2Ba4gydF4y2Ba- - - - - -gydF4y2Ba6gydF4y2Ba)输入数据在多个线程和分区进行密集模式匹配操作并行而引用一个DFA。然而,这种方法会导致大量的缓存错过DFA访问数量增加的模式字符串。我们的并行化方法,相反,分区给定的一组模式字符串分成多个组较小数量的模式字符串以节省空间的方式。因此,多个小DFA构造而不是单一的大型DFA的方法。使用多个小dfa,密集模式匹配操作同时进行的常见输入文本字符串。这导致显著较小的缓存的足印在每个核心的缓存引用分区dfa不规则的访问模式。因此,它导致较低的缓存错过比率和令人印象深刻的性能改进。实验结果在Nvidia Tesla GPU甘蓝型基于开普勒GK110体系结构表明,我们的方法导致了2.73倍的加速比以前输入数据分区的方法。吞吐量性能达到693 Gbps。相比单一CPU核心(6-core 2.0 Ghz Intel Xeon e5 - 2650),我们获得了加速在127 ~ 311的范围。 The speedup over the parallelized CPU version using 6 threads is in the range of 86~183. Experimental results on the Intel Xeon Phi with 61 x-86 cores also show up to 2.00 times speedup compared with the previous input data partitioning approach.

剩下的纸是组织如下。部分gydF4y2Ba2gydF4y2Ba描述了最新的许多核心加速器芯片的架构如Nvidia Tesla甘蓝型GPU和Intel Xeonφ。部分gydF4y2Ba3gydF4y2Ba介绍了AC算法。部分gydF4y2Ba4gydF4y2Ba描述我们的并行化方法,提高了缓存访问DFA的位置。部分gydF4y2Ba5gydF4y2Ba显示实验结果在Nvidia Tesla甘蓝型和Intel Xeonφ。部分gydF4y2Ba6gydF4y2Ba解释了之前的相关研究在并行AC算法。部分gydF4y2Ba7gydF4y2Ba概括了论文的结论。gydF4y2Ba

2。最近的许多核心加速器芯片架构gydF4y2Ba

最近,许多核心加速器芯片为HPC应用程序变得越来越流行。代表芯片是基于开普勒的Nvidia Tesla甘蓝型GK110架构和Intel Xeonφ基于许多集成的核心(MIC)体系结构。在接下来的部分,我们将描述这些架构。gydF4y2Ba

2.1。Nvidia Tesla GPU甘蓝型gydF4y2Ba

最新的GPU的体系结构的特点是大量的统一的精密可编程核或线程处理器为着色器取代了单独的处理单元,顶点和像素在前面的GPU。此外,最新的GPU的时钟频率显著增加。这些都大大提高了浮点gpu的性能,远远超过了最新的cpu。精密芯(或线程处理器)分布在多个流多处理器(SMX)(或线程块)(见图gydF4y2Ba1gydF4y2Ba)。软件线程分为许多线程组(称为扭曲)每个由32个线程。线程在相同的变形在计划和执行的线程处理器在同一SMX的SIMD(单指令多数据)模式。每个线程执行相同的指令由普通指令单元的数据流从缓存设备内存芯片上的记忆和寄存器。当运行经遇到一个缓存小姐,例如,上下文是转向了一种新扭曲而缓存小姐服务在接下来的几百个周期。因此,GPU多线程的方式执行。gydF4y2Ba

GPU是建立在一个复杂的内存层次结构如图gydF4y2Ba1gydF4y2Ba。有寄存器和当地的记忆属于每个线程处理器或核心。本地内存芯片外是一个区域设备内存。共享内存,一级缓存(L1)和只读数据缓存集成在一个线程阻塞的GPU。共享内存是一个快速(寄存器)一样快programmer-managed记忆。2级(L2)缓存集成GPU芯片和所有的线程块之间使用。全局内存是一个地区在芯片外设备内存访问的线程块,通过该技术可以与主机CPU通信。全局内存中的数据被程序员直接在共享内存缓存或者他们可以通过L1和L2缓存的缓存自动获得访问。有恒定的记忆和纹理内存区域在设备内存。这些区域是只读的数据。 They can be cached in the L2 cache and the read-only data cache. On Nvidia Tesla K20, the read-only data from the global memory can be loaded through the same cache used by the texture pipeline via a standard pointer without the need to bind to a texture beforehand. This read-only cache is used automatically by the compiler as long as certain conditions are met. The __限制gydF4y2Ba__预选赛时应该使用一个变量声明帮助编译器检测的条件(gydF4y2Ba10gydF4y2Ba]。gydF4y2Ba

为了有效地利用最新先进的GPU的体系结构、编程环境如CUDA (gydF4y2Ba10gydF4y2Ba从英伟达,OpenCL [gydF4y2Ba11gydF4y2Ba从Khronos组织],OpenACC [gydF4y2Ba12gydF4y2Ba从子群的OpenMP架构评审委员会(ARB)已经开发出来。使用这些环境,用户可以更直接控制大量的GPU核心及其复杂的内存层次结构。灵活的架构和编程环境导致了许多创新在许多应用领域的性能改进和更多的仍然是。gydF4y2Ba

2.2。Intel XeonφgydF4y2Ba

Intel Xeonφ(代号为骑士角落)是基于英特尔集成许多核心(MIC)架构相结合多个x86核心在单个芯片上。这种芯片可以运行在本机模式在应用程序运行时直接在它或在卸载模式应用程序运行在主机端,只有选中的区域(计算密集型地区)都被Xeonφ。卸载方式,Xeonφ是连接到一个主机通过pci - express Intel Xeon处理器总线。gydF4y2Ba

在本文中,我们使用的Xeonφ5120 d并行实现的交流:gydF4y2Ba(我)gydF4y2Ba这种协处理器60顺序计算核心支持64位x86指令。这些核心由高性能双向环互连连接(见图gydF4y2Ba2gydF4y2Ba)。它也有一个服务核心,因此总共有61核心芯片。gydF4y2Ba(2)gydF4y2Ba每个核心时速为1053 Mhz,提供了四通多线程(超),512位宽的SIMD向量对应八个双精度或16个单精度浮点数。gydF4y2Ba(3)gydF4y2Ba每个核心L1 32 KB数据缓存,32 KB L1指令缓存和512 KB统一L2高速缓存。因此,60核心结合30 MB L2高速缓存。L2高速缓存使用硬件支持是完全一致的。gydF4y2Ba(iv)gydF4y2BaXeonφ芯片有16个记忆频道提供5吨/ s(每秒Gigatransfer)传输速度。车载系统内存的总大小是8 GB。gydF4y2Ba

程序员可以使用相同的编程语言和模型的XeonφIntel Xeon处理器。它可以用Fortran编写的应用程序运行,C / c++,等等和并行模型如OpenMP, MPI, Pthreads,英特尔Clik +,和英特尔线程构建块(gydF4y2Ba13gydF4y2Ba]。gydF4y2Ba

3所示。Aho-Corasick (AC)字符串匹配算法gydF4y2Ba

Aho-Corasick (AC)是一种多模式字符串匹配算法可以同时匹配多个模式对于一个给定的有限集合的字符串(或字典)。AC算法由两个阶段组成。在第一阶段,模式匹配的机器称为交流自动机(机)是由一个给定的有限集的模式字符串。在第二阶段,构建交流电机用于查找字符串的位置模式在给定输入文本gydF4y2Ba9gydF4y2Ba]。gydF4y2Ba

一旦交流构造自动机,它调用三个函数在执行模式匹配的第二阶段:goto函数gydF4y2Ba ,一个失败的功能gydF4y2Ba ,和一个输出函数gydF4y2Ba输出gydF4y2Ba。图gydF4y2Ba3gydF4y2Ba显示这些功能对于一个给定的模式gydF4y2Ba “他”、“她”、“他”、“她”}gydF4y2Ba9gydF4y2Ba]:gydF4y2Ba(我)gydF4y2Ba图的有向图gydF4y2Ba3(一个)gydF4y2Ba代表了goto函数gydF4y2Ba 。的gydF4y2Ba 函数映射一个国家和一个输入符号组成的一对成一个州或一个信息gydF4y2Ba失败gydF4y2Ba。例如,边缘标记gydF4y2BahgydF4y2Ba从0到1表明状态gydF4y2BaggydF4y2Ba(0,'gydF4y2BahgydF4y2Ba')= 1。(gydF4y2Ba ('gydF4y2BahgydF4y2Ba”、“gydF4y2Ba年代gydF4y2Ba”)表示所有输入符号以外的gydF4y2BahgydF4y2Ba”、“gydF4y2Ba年代gydF4y2Ba”)。没有一个箭头表示失败。交流电机具有的属性gydF4y2Ba 失败gydF4y2Ba对所有输入符号gydF4y2Ba 。gydF4y2Ba(2)gydF4y2Ba失败的函数gydF4y2Ba (图gydF4y2Ba3 (b)gydF4y2Ba)地图一个状态到另一个状态。这是咨询每当goto函数报告一个“失败”。gydF4y2Ba(3)gydF4y2Ba输出函数gydF4y2Ba输出gydF4y2Ba(图gydF4y2Ba3 (c)gydF4y2Ba)的一组关键字映射到输出到指定的状态。gydF4y2Ba

AC算法可以实现非确定性有限自动机(NFA)或确定性有限自动机(DFA)。在这篇文章中,我们实现了AC算法作为一个DFA的代表所有可能的状态机以及信息的可接受的状态转换系统(gydF4y2Ba14gydF4y2Ba]。DFA是由一组有限的状态gydF4y2Ba 和下一步行动的功能gydF4y2Ba 这样,每个状态gydF4y2Ba 和输入符号gydF4y2Ba ,gydF4y2Ba 是一个国家gydF4y2Ba 。因此,下一步的功能gydF4y2Ba 代替使用goto功能和故障函数中引入图gydF4y2Ba3gydF4y2Ba。输出函数也纳入DFA。图gydF4y2Ba4gydF4y2Ba显示了一组模式的交流电机gydF4y2Ba 他,她,他,和她},这三个功能gydF4y2Ba ,gydF4y2Ba ,gydF4y2Ba输出gydF4y2Ba集成在一个DFA。从初始状态开始,交流电机接受一个输入字符,并从当前状态转移到下一个正确的状态。gydF4y2Ba

伪代码gydF4y2Ba1gydF4y2Ba显示了交流电机是DFA。在这个伪代码,gydF4y2Ba 字符输入文本字符串的读取顺序执行循环。下一步功能gydF4y2Ba 从当前状态获取新的状态和角色gydF4y2Ba 。在新状态,检查是否存在匹配算法(如果gydF4y2Ba输出gydF4y2Ba(国家)! =空)。如果是这样的话,执行打印输出函数匹配的模式。在这段代码中,我们只使用两个函数:下一步的功能gydF4y2Ba 和输出功能。函数删除失败而NFA到DFA转换。假设我们有模式集gydF4y2Ba 他,她,他,她}和输入文本字符串”招待员。“DFA工作方式如下:gydF4y2Ba(我)gydF4y2Ba自gydF4y2Ba (0,'gydF4y2BaugydF4y2Ba)= 0,交流电机反馈状态0。gydF4y2Ba(2)gydF4y2Ba自gydF4y2Ba (0,'gydF4y2Ba年代gydF4y2Ba')= 3,gydF4y2Ba (3,gydF4y2BahgydF4y2Ba)= 4,gydF4y2Ba (4”,gydF4y2BaegydF4y2Ba')= 5,交流电机发出输出(5)=gydF4y2Ba 他,她}。gydF4y2Ba(3)gydF4y2Ba自gydF4y2Ba (5’gydF4y2BargydF4y2Ba')= 8gydF4y2Ba (8。”gydF4y2Ba年代gydF4y2Ba)= 9,交流电机发出输出(9)=gydF4y2Ba 。gydF4y2Ba

/gydF4y2Ba 输入:输入文本gydF4y2Ba ,gydF4y2Ba =输入文本的长度gydF4y2Ba
输出:关键词出现的位置gydF4y2Ba /gydF4y2Ba
过程DFA_AC(字符gydF4y2Ba ,intgydF4y2Ba )gydF4y2Ba
开始gydF4y2Ba
int状态= 0;gydF4y2Ba
for (intgydF4y2Ba ;gydF4y2Ba ;gydF4y2Ba + +)gydF4y2Ba
开始gydF4y2Ba
状态=gydF4y2Ba (状态,gydF4y2Ba );gydF4y2Ba
如果(gydF4y2Ba输出gydF4y2Ba(国家)! =空)gydF4y2Ba
开始gydF4y2Ba
打印gydF4y2Ba
打印gydF4y2Ba输出gydF4y2Ba(状态)gydF4y2Ba
结束gydF4y2Ba
结束gydF4y2Ba
结束gydF4y2Ba

AC算法的复杂性gydF4y2Ba ,在那里gydF4y2Ba 模式的长度之和,gydF4y2Ba 输入文本的长度,gydF4y2Ba 是模式的数量出现在输入文本。自动机的建设gydF4y2Ba 。基于自动机的模式匹配操作gydF4y2Ba 。提前知道组模式时,不改变在运行时如电脑病毒扫描,自动机的建设可以进行一次脱机和自动机多次用于模式匹配在第二阶段。在这种情况下,复杂性gydF4y2Ba 。gydF4y2Ba

4所示。缓存Locality-Centric并行化gydF4y2Ba

在本节中,我们解释我们的并行化方法。我们首先描述初步的步骤。然后,我们描述我们的方法,分区DFA分成多个小块为提高缓存位置。gydF4y2Ba

4.1。预赛gydF4y2Ba
以下4.4.1。AC算法的执行情况gydF4y2Ba

节中解释gydF4y2Ba3gydF4y2Ba、交流模式匹配DFA对于一个给定的有限集合的字符串(或字典)是构建在第一阶段。我们假设构造DFA是固定在第二阶段的AC DFA用于查找字符串的位置在给定的输入文本模式。这样一个很好的例子就是病毒的杀毒软件数据库(DFA)是由一组给定的几十名已知病毒(gydF4y2Ba1gydF4y2Ba在第一阶段。然后,密集模式匹配操作进行检测硬盘中的病毒图像使用病毒数据库在第二阶段。后者阶段耗时和多次重复使用相同的病毒数据库之前,用户更新它。因此,在本文中,我们假设AC DFA构造一旦在单CPU主机处理器的核心在第一阶段,是在GPU上进行并行字符串匹配使用DFA在第二阶段我们的并行化。gydF4y2Ba

4.1.2。构造DFAgydF4y2Ba

交流模式匹配DFA在第一阶段被构建为一个二维矩阵称为状态转换表(STT)(见图gydF4y2Ba5gydF4y2Ba)。行STT代表美国的DFA和列表示输入字符。因此,对于一个给定的状态gydF4y2Ba 和一个输入字符gydF4y2Ba ,一个条目STT [gydF4y2Ba ][gydF4y2Ba ]表示相应的下一个状态或故障状态。假设我们有256个输入字符(映射到256个字符的扩展ASCII表),然后是STT需要257列:256列人物和1列指示如果当前状态是一个匹配的状态,输出函数的执行打印位置和状态的输入数据中发现模式。gydF4y2Ba

4.1.3。数据位置gydF4y2Ba

一旦STT由主机CPU和存储在主机内存,我们将其复制到设备内存的GPU的输入数据。当复制这些数据,我们需要仔细地决定在设备内存的GPU(全局内存,不变的记忆,或纹理内存)我们需要存储它们。生成大量的数据访问输入数据和STT数据而进行模式匹配操作。从一开始输入数据访问顺序。另一方面,STT显示不同的访问模式。从DFA的初始状态开始,一个字符的输入数据查找当前状态找到下一个状态。在编译时不知道这状态转换信息;因此,它会导致STT接近随机数据访问模式。gydF4y2Ba

考虑到上述数据访问的特点,我们把输入数据的全局内存,以便它可以自动加载到L1缓存(通过片上L2高速缓存)或显式地加载到共享内存GPU的程序员访问延迟最小化。我们把STT纹理内存中的数据,以便积极使用部分随机访问期间可以自动缓存的L2和只读缓存GPU。这种分离的访问路径输入和STT以便他们不直接相互干扰。因此,它减少了内存访问延迟和更有效地使用可用内存带宽。图gydF4y2Ba6gydF4y2Ba显示了生成的数据访问产生的多个线程分配给多个线程的处理器(或精密核)上的线程块GPU执行模式匹配操作时根据我们的数据布局方案。gydF4y2Ba

4.1.4。其他注意事项有效的并行化gydF4y2Ba

输入数据和STT数据存储在全局内存和纹理内存,分别,我们考虑下面的一个高效的并行化和高性能模式匹配操作:gydF4y2Ba(我)gydF4y2Ba如上所述gydF4y2Ba4.1.3gydF4y2Ba,每个输入数据块分配给一个线程在线程阻塞。当进行模式匹配操作,我们跨越每个线程的输入块gydF4y2Ba 字符,gydF4y2Ba 的最大模式长度字符串的模式。通过这样做,当我们避免错过一个模式匹配一个模式字符串位于两个相邻的边界数据块有两个不同的线程。gydF4y2Ba(2)gydF4y2Ba相同的硬件上运行的软件线程阻塞线程阻塞(在经单位)生成一个输入数据的访问和STT。随着GPU多线程的方式执行,长期全局内存访问延迟输入数据或经STT数据可以掩盖掉或被其他扭曲的模式匹配操作属于同一个线程块或其他线程块。gydF4y2Ba(3)gydF4y2Ba全球加载内存访问开销可以进一步减少输入数据通过有效地利用芯片上的共享内存。我们首先将输入数据划分为若干块。每个数据块分配给每个线程块。所有线程在一块合作加载相应的数据块从全局内存共享内存执行模式匹配操作。模式匹配操作一块线程执行多线程的方式变形单元的输入数据加载和STT全局内存的数据加载其他的扭曲。为了使用共享内存以最优的方式,我们需要仔细确定线程块的数量和每个线程阻塞的线程数量。我们将更详细地解释这部分gydF4y2Ba4.3gydF4y2Ba。gydF4y2Ba(iv)gydF4y2Ba同时将数据加载到共享内存,一个重要的考虑是联合全球内存访问性能。在我们的并行实现,让多个线程块配合加载一个又一个的数据块完全加载数据块的线程阻塞。我们将更详细地描述我们的全球内存访问合并技术在部分gydF4y2Ba4.3gydF4y2Ba。gydF4y2Ba(v)gydF4y2Ba当输入数据块从全局内存加载到共享内存,我们需要仔细的存储方案来避免或者减少共享内存存储数据时银行冲突被多个线程同时访问也分布在多个银行。我们使用一个存储方案的输入数据从全局内存加载分割成小的字节数和存储不同的银行,以避免任何银行冲突。我们将更详细地描述我们计划部分gydF4y2Ba4.3gydF4y2Ba。gydF4y2Ba

4.2。基于DFA分区并行化gydF4y2Ba

一旦输入数据和STT数据放置在全球内存和纹理内存密集型模式匹配的操作是在第二阶段进行交流时引用这些数据。在前面的研究(gydF4y2Ba4gydF4y2Ba- - - - - -gydF4y2Ba6gydF4y2Ba),他们通过分区并行第二阶段输入数据到多个块,每一块分配给不同的处理器核心或线程。然后,每个核心或线程进行模式匹配在引用单个大STT并行操作。这种方法,然而,海量缓存错过开销为访问STT STT访问很不规则。此外,随着模式字符串数量的增加,规模的STT相应增加。因此,一个大的STT随机访问并行多核或线程导致穷人缓存位置和较低的性能。gydF4y2Ba

为了显著降低管理费用与不规则STT访问高缓存错过,我们给定的一组模式分割成多个小模式集。然后,每个小模式组,我们构建一个相应的DFA在第一阶段的交流表示为多个小stt(见图gydF4y2Ba7 (b)gydF4y2Ba)。在第二阶段,整个输入数据加载多个核心或线程使用同一个STT模式匹配。图gydF4y2Ba8gydF4y2Ba比较我们的并行化方法(图gydF4y2Ba8 (b)gydF4y2Ba(图)与前面的方法gydF4y2Ba8(一个)gydF4y2Ba)。以前的方法(gydF4y2Ba4gydF4y2Ba- - - - - -gydF4y2Ba6gydF4y2Ba)使用分割输入数据在多个核心或线程和普通大STT。我们的方法使用分区小stt (gydF4y2Ba )从多核或线程和常见的输入数据。因此,我们的方法可以显著减少缓存足迹引用每个核心的STT或线程。我们的方法,另一方面,整个输入数据为每个核心或线程。由于输入数据按顺序访问,它可以有效地从全局内存加载到芯片上的共享记忆GPU的程序员。因此,我们的方法会导致更好的缓存命中率和整体性能,稍后我们将展示部分gydF4y2Ba5gydF4y2Ba(实验结果)。gydF4y2Ba

当我们给定的一组模式分割成多个小模式集,我们使用一个算法包括两个部分:gydF4y2Ba(我)gydF4y2Ba第1部分(算法gydF4y2Ba1gydF4y2Ba)分配模式字符串相同的字符到一个STT开始。它分配模式循环的方式从模式集最多的次数最少的出现。在5 ~ 6行,我们计算模式的数量的字符开始gydF4y2Ba 。这一步形成了gydF4y2Ba 组包含256个元素对应的256个字符gydF4y2Ba 集。然后,gydF4y2Ba 设置和gydF4y2Ba 设置按降序排列(7和8行)。通过这种安排,大量出现的人物被放置到第一的位置gydF4y2Ba 集。从9 ~ 13行,该算法计算的位置模式分配给组。这些职位都是计算使用循环分布。在第1部分帮助分发STT之间大约相同数量的模式,可能会有一些差异产生的STT大小。gydF4y2Ba(2)gydF4y2Ba第2部分(算法gydF4y2Ba2gydF4y2Ba)平衡模式的数量在每个STT STT之间通过重新分配模式。的模式是与最多的模式从STT的STT最少的模式。嵌套使用while循环(5 ~ 11行)使这种转变。当进入内部while循环,我们检查是否两套长度差异(gydF4y2Ba ,gydF4y2Ba )最大长度和最小长度比一个大gydF4y2Ba 值(gydF4y2Ba )。如果是这样,我们将从一个模式gydF4y2Ba 来gydF4y2Ba 和更新gydF4y2Ba 和gydF4y2Ba 。我们重复这个步骤,直到(gydF4y2Ba )等于或小于gydF4y2Ba 。后内gydF4y2Ba 循环退出时,位置和更新所有的模式集的总长度(9 ~ 10行)。的gydF4y2Ba 值是在外部使用gydF4y2Ba 循环检查stt的大小差异。如果大小的区别gydF4y2Ba 和gydF4y2Ba 大于gydF4y2Ba 我们进入了内循环,平衡gydF4y2Ba 和gydF4y2Ba 。的gydF4y2Ba 值是选择后进行大量的实验。我们设置阈值= 15%的总数小于20000模式,模式的数量大于20000 10%,分别。因此,使用我们的算法,所有由此产生的stt不同大小不大于15%的< 20000 >模式和10%的20000年模式。gydF4y2Ba

输入gydF4y2Ba:gydF4y2Ba ,在那里,gydF4y2Ba
:数量的模式字符串gydF4y2Ba
:stt我们想生成的数量(gydF4y2Ba 2的倍数)gydF4y2Ba
,一组模式字符串gydF4y2Ba
输出gydF4y2Ba:gydF4y2Ba ,在那里,gydF4y2Ba
和gydF4y2Ba 包括属于模式字符串gydF4y2Ba
(gydF4y2Ba )给gydF4y2Ba ,一组256个字符的ASCII表gydF4y2Ba
(gydF4y2Ba )让gydF4y2Ba
(gydF4y2Ba )(我)gydF4y2Ba 模式字符串的数量gydF4y2Ba 开始与性格gydF4y2Ba
(gydF4y2Ba )(二)gydF4y2Ba 和gydF4y2Ba
(gydF4y2Ba )gydF4y2BaforeachgydF4y2Ba 做gydF4y2Ba
(gydF4y2Ba )计算gydF4y2Ba
(gydF4y2Ba )排序gydF4y2Ba 在降序排列gydF4y2Ba
(gydF4y2Ba )安排的位置gydF4y2Ba 基于的顺序gydF4y2Ba
(gydF4y2Ba )gydF4y2BaforeachgydF4y2Ba 开始与性格gydF4y2Ba 做gydF4y2Ba
(gydF4y2Ba )gydF4y2Ba如果gydF4y2Ba 然后gydF4y2Ba
(gydF4y2Ba )分配gydF4y2Ba 来gydF4y2Ba ,在那里gydF4y2Ba
(gydF4y2Ba )gydF4y2Ba其他的gydF4y2Ba
(gydF4y2Ba )分配gydF4y2Ba 来gydF4y2Ba ,在那里gydF4y2Ba

输入gydF4y2Ba:gydF4y2Ba ,在那里,gydF4y2Ba
和gydF4y2Ba 包括属于模式字符串gydF4y2Ba
输出gydF4y2Ba:gydF4y2Ba ,在那里,gydF4y2Ba
和gydF4y2Ba 平衡模式的长度吗gydF4y2Ba
(gydF4y2Ba )让gydF4y2Ba
(gydF4y2Ba )(我)gydF4y2Ba 模式的总长度gydF4y2Ba ,在那里gydF4y2Ba
(gydF4y2Ba )(二)gydF4y2Ba 和gydF4y2Ba has-maximum-length的位置gydF4y2Ba 和has-minimum-lengthgydF4y2Ba 分别gydF4y2Ba
(gydF4y2Ba )(3)gydF4y2Ba 是限制的长度两个stt之间的区别gydF4y2Ba
(gydF4y2Ba )gydF4y2Ba做gydF4y2Ba
(gydF4y2Ba )gydF4y2Ba而gydF4y2Ba 做gydF4y2Ba
(gydF4y2Ba )移动一个模式在最后的位置gydF4y2Ba 来gydF4y2Ba
(gydF4y2Ba )更新gydF4y2Ba 和gydF4y2Ba
(gydF4y2Ba )更新模式集的总长度gydF4y2Ba
(gydF4y2Ba )更新gydF4y2Ba 和gydF4y2Ba
(gydF4y2Ba )gydF4y2Ba而gydF4y2Ba ;gydF4y2Ba

最优分布会产生各stt相同数量的模式,所有stt的大小相等。所有的组合和STT最小化(小至或接近原单大STT生成的大小在前面的方法)。使用我们的算法,生成的多个DFA的组合规模匹配的大小单大DFA中我们将展示部分gydF4y2Ba5。4gydF4y2Ba。因此,我们的方法构造多个dfa空间。此外,它花费更少的时间来构造多个小dfa与时间构建一个大的STT使用前面的方法,我们还将展示部分gydF4y2Ba5。4gydF4y2Ba。gydF4y2Ba

DFA分区算法时间效率,因为线性两种算法的复杂性gydF4y2Ba1gydF4y2Ba和gydF4y2Ba2gydF4y2Ba。让我们假设如下:gydF4y2Ba(1)gydF4y2Ba :模式的数量;gydF4y2Ba(2)gydF4y2Ba :零件的数量(dfa或stt)我们想分区。gydF4y2Ba算法gydF4y2Ba1gydF4y2Ba由四个任务:gydF4y2Ba(我)gydF4y2Ba计算值gydF4y2Ba (5 ~ 6行),gydF4y2Ba(2)gydF4y2Ba在设置值gydF4y2Ba (第7行),gydF4y2Ba(3)gydF4y2Ba排序值gydF4y2Ba (第8行),gydF4y2Ba(iv)gydF4y2Ba分配模式集(9 ~ 13行)。gydF4y2Ba

我们可以计算任务的执行时间如下:gydF4y2Ba(1)gydF4y2Ba(2)和(3)总是常数gydF4y2Ba (事实上他们的时间复杂性gydF4y2Ba 。)gydF4y2Ba(2)gydF4y2Ba(我),(四)=gydF4y2Ba ,从而gydF4y2Ba 。gydF4y2Ba因此,算法的复杂性gydF4y2Ba1gydF4y2Ba是gydF4y2Ba 。gydF4y2Ba

算法后gydF4y2Ba1gydF4y2Ba执行,我们的模式划分为gydF4y2Ba 部分。因此,gydF4y2Ba(1)gydF4y2Ba每个部分都有gydF4y2Ba 模式在最好的情况下,gydF4y2Ba或gydF4y2Ba(2)gydF4y2Ba一组接近gydF4y2Ba 模式和其他为最坏的情况下几乎是空的。gydF4y2Ba算法gydF4y2Ba2gydF4y2Ba有两个循环。在最坏的情况下,内循环执行gydF4y2Ba 次,因为一个STT接近gydF4y2Ba 模式和其他STT几乎是空的。外层循环,我们需要更新的长度设置并选择两组最大长度和最小长度为下一步。这个过程选择两组gydF4y2Ba 集。因此,执行时间成正比gydF4y2Ba 。总执行时间=gydF4y2Ba 。因此,算法的复杂性gydF4y2Ba2gydF4y2Ba是gydF4y2Ba ,因为gydF4y2Ba 是一个常数(stt)。gydF4y2Ba

4.3。进一步的性能优化gydF4y2Ba

除了上述基于DFA分区并行化的方法,我们将进一步性能优化技术应用到GPU实现。他们大多是来自我们的早期工作gydF4y2Ba15gydF4y2Ba]。gydF4y2Ba(我)gydF4y2Ba输入数据存储在一个连续的方式在全球记忆。加载数据块时,我们让多个线程生成内存访问属于128字节的范围,这些访问合并为一个请求和发送到内存(gydF4y2Ba10gydF4y2Ba]。这节省了内存带宽,提高了性能(gydF4y2Ba15gydF4y2Ba]。gydF4y2Ba(2)gydF4y2Ba输入数据后从全局内存加载,每个线程访问存储在共享内存块的输入。当多个线程试图同时读取自己的数据块分布在多个银行的共享内存,它将导致很多银行冲突。我们使用一个存储方案通过分配一个数据块从全局内存加载被划分成4个字节单元和存储在共享内存的地址连续映射到共享内存银行对角线。这种存储方案避免了任何银行冲突和结果在一个无冲突负载从共享内存的银行(gydF4y2Ba15gydF4y2Ba]。gydF4y2Ba(3)gydF4y2BaGPU以多线程的方式执行。理论上有多个线程同时执行可以容忍片外存储器(全局内存、纹理内存,等等)的访问延迟以几百周期。然而,片外存储器的带宽限制。如果有太多的并发访问片外存储器,它可以导致内存访问路径和拥堵进一步延长延迟。此外,导致增加缓存的线程数量增加了(gydF4y2Ba16gydF4y2Ba]。因此,找到一个最优数量的线程来有效地隐藏了片外存储器延时而有效地利用大量的线程处理器和内存带宽获得高绩效是至关重要的。我们试图找到并安排最佳的并行线程数到硬件线程块和线程处理器核心近详尽的搜索各种输入块大小是分配给每个线程(gydF4y2Ba15gydF4y2Ba]。gydF4y2Ba我们将展示上述优化技术的性能好处除了我们多个基于stt后部分gydF4y2Ba5.2。2gydF4y2Ba。gydF4y2Ba

5。实验结果gydF4y2Ba

在本节中,我们首先解释我们的实现细节基于DFA分区并行AC算法。然后,我们目前的实验结果对特斯拉甘蓝型GPU和Intel Xeonφ。为了证明我们的方法的空间效率,我们还比较单一的大小STT(以前的方法)与多个较小的STT之和(我们的方法)。我们也显示建筑的成本比较多个STT与一个大的STT相比,我们的方法在前面的方法。因此,我们证明了我们的方法的时间效率建设STT也。gydF4y2Ba

5.1。实现细节gydF4y2Ba

我们的实验系统上进行了合并主机Intel Xeon多核处理器(2.0 Ghz Intel Xeon e5 6-core - 2650)与20 MB三级缓存,5 GB的Nvidia Tesla甘蓝型GPU设备内存,和Intel Xeonφ61 x - 86核与8 GB的内存。我们也使用Linux Centos 5.5。下面我们描述的方法来生成测试输入数据和模式的数据。我们也解释了有关并行实现。gydF4y2Ba

5.1.1。测试数据生成gydF4y2Ba

为了生成随机输入数据集和参考模式数据集用于我们的实验中,我们首先收集了50 GB的数据从各种各样的英语杂志,如时间和英国广播公司(BBC),和其他很多。然后,我们随机输入数据和模式中提取数据从收集到的数据。我们使用的输入数据大小的范围20 MB ~ 500 MB。使用的模式数量在100 ~ 50000的范围。我们也生成一个特殊的输入数据,gydF4y2BaDict_InputgydF4y2Ba。有两种gydF4y2BaDict_InputgydF4y2Ba:(1)gydF4y2BaDict_Input_SgydF4y2Ba,输入数据生成的内容直接从字典中所有模式字符串。因此,gydF4y2BaDict_Input_SgydF4y2Ba有体积小。例如,当模式在字典里的数量是100(50000)和模式长度是10个字符,平均的大小gydF4y2BaDict_Input_SgydF4y2Ba大约是1 KB (512 KB)。(2)在gydF4y2BaDict_Input_LgydF4y2Ba输入生成的内容,复制和连接所有模式在字典里输入规模大。输入数据总结在表的信息gydF4y2Ba1gydF4y2Ba。模式集的特点给出了表gydF4y2Ba2gydF4y2Ba。gydF4y2Ba

5.1.2中。并行实现gydF4y2Ba

特斯拉甘蓝型GPU实现,我们使用共享内存加载输入文本数据。我们还展示了实现不使用共享内存,以量化的好处使用共享内存。我们描述的实现如下:gydF4y2Ba(我)gydF4y2Bap - 1:全局内存只(或没有共享内存)实现(参见图gydF4y2Ba9gydF4y2Ba)输入文本数据复制到全局内存。然后,积极使用部分的输入数据缓存到芯片上的自动缓存(L2和L1缓存中),但它不是显式地在共享内存缓存。STT数据附加到纹理内存和使用的积极部分STT数据缓存的L2缓存和只读数据缓存。在这个实现中,使用的L2缓存输入数据和STT数据。因此,缓存的性能影响locality-centric并行化方法更杰出的有效L2高速缓存大小使用的输入数据和STT数据得到减少。gydF4y2Ba(2)gydF4y2Bap 2:共享内存实现(参见图gydF4y2Ba10gydF4y2Ba)将输入数据的全局内存加载到芯片上的共享记忆由程序员显式。STT数据放置在纹理内存和使用的积极部分加载到L2高速缓存的方式,然后在只读数据缓存中。因此,L2缓存STT使用数据。在这个实现中,输入数据缓存与p - 1相比是更有效的实现依赖于自动缓存。STT数据缓存也变得更有效率,因为L2高速缓存现在致力于STT数据。gydF4y2Ba

为了实现交流的多个基于stt字符串匹配算法,我们使用CUDA流特性。在费米架构与流特性,只有16路支持并发和溪流多路复用到一个队列,开普勒甘蓝型允许32位并发性和一个工作队列每流。这导致了所有信息流级别的并发性,没有inner-stream依赖(gydF4y2Ba17gydF4y2Ba]。因此,我们创建多个CUDA流的数量等于STT(4或8流在我们的例子中),并将每个流分配给每个模式匹配任务在一个较小的STT引用可能匹配与整个输入数据。这保证了模式匹配可以同时使用多个stt执行任务。为了存储stt纹理内存,我们使用一个新特性称为纹理对象(或bindless纹理,因为他们不需要手动绑定/释放)从开普勒架构(CUDA 5.0或更高版本)。创建纹理对象的数量等于stt的数量。我们只将这些纹理对象传递给内核使用。gydF4y2Ba

伪代码gydF4y2Ba2gydF4y2Ba显示在甘蓝型GPU的伪代码实现。首先,创建的纹理对象绑定到stt(2 ~ 7行)。接下来,我们创建一个流(9 ~ 14行)。然后,流配合将输入数据复制到设备内存。(每个流拷贝一个数据段(16 ~ 19行)。)输入数据复制后,每个流执行模式匹配任务使用其输入数据和STT数据(21 ~ 23行)。最后,结果复制回主机,然后流被破坏(25 ~ 31行)。gydF4y2Ba

(gydF4y2Ba )gydF4y2Ba⋯gydF4y2Ba
(gydF4y2Ba )gydF4y2BacudaArray * cuda_arrays =gydF4y2Bamalloc (cudaArray * *)gydF4y2Ba(num_of_sttsgydF4y2Ba*gydF4y2Basizeof (cudaArray *))gydF4y2Ba
(gydF4y2Ba )gydF4y2BacudaTextureObject_t * textObj = (cudaTextureObject_t *)gydF4y2Bamalloc (num_of_stts *运算符gydF4y2Ba
(cudaTextureObject_t));gydF4y2Ba
(gydF4y2Ba )gydF4y2Ba为gydF4y2Ba(int i = 0;我< num_of_stts;gydF4y2Ba我+ +)gydF4y2Ba
(gydF4y2Ba )gydF4y2Bacuda_arrays[我]= generate_cuda_arraygydF4y2Ba(get_dfa_matrix(我),gydF4y2Baget_dfa_width(我),gydF4y2Ba
get_dfa_height (i));gydF4y2Ba
(gydF4y2Ba )gydF4y2BatextObj[我]= generate_texture_objectsgydF4y2Ba(cuda_arrays[我]);gydF4y2Ba
(gydF4y2Ba )gydF4y2Ba
(gydF4y2Ba )gydF4y2Ba⋯gydF4y2Ba
(gydF4y2Ba )gydF4y2BacudaStream_tgydF4y2Ba*流=gydF4y2Ba
(gydF4y2Ba )gydF4y2Ba(cudaStream_t *) mallocgydF4y2Ba(nstreams * sizeof (cudaStream_t));gydF4y2Ba
(gydF4y2Ba )gydF4y2Ba 创建多个CUDA流gydF4y2Ba
(gydF4y2Ba )gydF4y2Ba为gydF4y2Ba(int i = 0;我< nstreams;gydF4y2Ba我+ +)gydF4y2Ba
(gydF4y2Ba )gydF4y2BacudaStreamCreategydF4y2Ba[我])(&(流);gydF4y2Ba
(gydF4y2Ba )gydF4y2Ba
(gydF4y2Ba )gydF4y2Ba/ /将数据复制到GPU内存,每个流拷贝一个段gydF4y2Ba
(gydF4y2Ba )gydF4y2Ba为gydF4y2Ba(int i = 0;我< nstreams;我+ +)gydF4y2Ba
(gydF4y2Ba )gydF4y2Ba长in_offset =我* input_len / nstreams;gydF4y2Ba
(gydF4y2Ba )gydF4y2BacudaMemcpyAsyncgydF4y2Ba(d_input + in_offsetgydF4y2Bah_input + in_offset,gydF4y2Bainput_len *运算符gydF4y2Ba
(字符)/ nstreams,gydF4y2BacudaMemcpyHostToDevice,gydF4y2Ba流[我]);gydF4y2Ba
(gydF4y2Ba )gydF4y2Ba
(gydF4y2Ba )gydF4y2Ba 每个流处理输入数据与每个dfa (texObj[我])gydF4y2Ba
(gydF4y2Ba )gydF4y2Bafor (int i = 0;我< nstreams;gydF4y2Ba我+ +)gydF4y2Ba
(gydF4y2Ba )gydF4y2Ba匹配< < <块,gydF4y2Ba线程,gydF4y2Basm_size,gydF4y2Ba流[我]> > > (texObj[我],gydF4y2Bad_input,gydF4y2Ba
input_len,gydF4y2Bapattern_max_len,gydF4y2Bad_output,gydF4y2Baoutput_len);gydF4y2Ba
(gydF4y2Ba )gydF4y2Ba
(gydF4y2Ba )gydF4y2Ba 复制结果返回给主机CPU,每个流拷贝一个段gydF4y2Ba
(gydF4y2Ba )gydF4y2Ba为gydF4y2Ba(int i = 0;我< nstreams;gydF4y2Ba我+ +)gydF4y2Ba
(gydF4y2Ba )gydF4y2Ba长out_offset =我* output_len / nstreams;gydF4y2Ba
(gydF4y2Ba )gydF4y2BacudaMemcpyAsyncgydF4y2Ba(h_output + out_offsetgydF4y2Bad_output + out_offset,gydF4y2Baoutput_len *运算符gydF4y2Ba
(int) / nstreams,gydF4y2BacudaMemcpyDeviceToHost);gydF4y2Ba
(gydF4y2Ba )gydF4y2Ba
(gydF4y2Ba )gydF4y2Ba为gydF4y2Ba(int i = 0;我< nstreams;gydF4y2Ba我+ +)gydF4y2Ba
(gydF4y2Ba )gydF4y2BacudaStreamDestroygydF4y2Ba(流[我]);gydF4y2Ba
(gydF4y2Ba )gydF4y2Ba
(gydF4y2Ba )gydF4y2Ba⋯gydF4y2Ba

5.2。在甘蓝型GPU性能比较gydF4y2Ba

我们将展示我们的方法与以前的方法的性能结果。在所有的实验中,我们显示的时间进行模式匹配操作只是因为AC算法并行化的第二阶段。gydF4y2Ba

5.2.1。我们的方法在过去的方法的性能优势gydF4y2Ba

图gydF4y2Ba11gydF4y2Ba显示了p - 1的吞吐量性能(仅全局内存)实现一系列的输入数据大小(20 ~ 500 MB)和一系列的数字模式(100、5000和50000年)以Gbps。性能比较,我们还实现了之前的方法,单一大STT和分区输入文本数据块用于进行并行模式匹配。(图中标注1 STT显示了前面的方法的性能)。gydF4y2Ba(我)gydF4y2Ba我们的新方法(使用4、8 STT)优于之前的方法,使用单一的大型STT。随着输入数据大小的增加,我们的方法的性能稳步提高到100 MB,然后开始饱和。在p - 1方法中,L2高速缓存由输入数据和共享STT;因此,随着输入数据大小的增加,压强的L2缓存相应增加,导致性能饱和度。然而,我们的方法和前面的方法之间的性能差距被拉大。(稍后我们将展示这一点。p 2的实验结果,性能饱和度增加输入数据大小消失的L2缓存stt使用。)gydF4y2Ba(2)gydF4y2Ba模式的数量增加时,吞吐量性能得到较低的在所有的情况下因为缓存未命中增加数量增加的模式。更多的模式影响前面的方法更加的性能。因此,性能差距正在扩大。gydF4y2Ba

图gydF4y2Ba12gydF4y2Ba显示p 2的吞吐量性能(使用共享内存)实现一系列的输入数据大小(50 ~ 500 MB)和一系列的数字模式(100、5000和50000年)以Gbps。gydF4y2Ba(我)gydF4y2Ba在p - 1的结果,这种新方法(使用4、8 STT)优于之前的方法(1 STT)。随着输入数据大小的增加,性能差距变大。我们的方法提高了性能进一步提高输入数据的大小500 MB。gydF4y2Ba(2)gydF4y2Ba随着模式的数量增加,吞吐量性能得到较低的在所有的情况下。然而,我们的方法和单一STT之间的性能差距变大如p - 1实现的方法。gydF4y2Ba(3)gydF4y2Ba最好的表现为p 2 p - 1实现75.8 Gbps, 692.7 Gbps。因此,共享内存实现放弃~ 9.14倍比p - 1实现更好的性能。gydF4y2Ba

图gydF4y2Ba13gydF4y2Ba展示了我们的方法在过去的加速方法。gydF4y2Ba(我)gydF4y2Ba图gydF4y2Ba(13日)gydF4y2Ba显示,使用p - 1实现,我们的方法给出了加速在1.47 ~ 2.73的范围从20 MB到500 MB的数据大小和数量的模式从100年到50000年不等。随着输入数据的大小增加,加速提高。模式的数量直接对所有输入性能影响大小为500 MB。gydF4y2Ba(2)gydF4y2Ba图gydF4y2Ba13 (b)gydF4y2Ba展示了我们的方法的加速使用共享内存(p 2)实现。我们的方法的结果在1.34 ~ 1.86的加速数据大小从20 MB到500 MB和模式的数量从100年到50000年不等。随着数据量的增加,加速了增加到200 MB。除了200 MB,饱和烃的加速增长。随着模式的数量增加,加速相应增加。gydF4y2Ba(3)gydF4y2Ba总体加速p 2较低,然而,比p - 1的加速。在使用共享内存的p 2, L2高速缓存致力于STT访问。因此,它可以捕获更大的部分的工作集大型单STT p 2中使用前面的方法实现。当模式的数量增加到50000,但是,我们看到一个突然增加的加速。工作集的大单STT使用在前面的方法开始溢出L2缓存。这表明,我们的方法的有效性,为更多的更大的模式。gydF4y2Ba

图gydF4y2Ba(14日)gydF4y2Ba显示了加速的结果p 2实现在单CPU核心(6-core 2.0 Ghz Intel Xeon e5 - 2650)。加速范围在127 ~ 311。加速实现p 2的六线平行版本范围在86 ~ 183,如图gydF4y2Ba14 (b)gydF4y2Ba。gydF4y2Ba

数据gydF4y2Ba(15日)gydF4y2Ba和gydF4y2Ba15 (b)gydF4y2Ba的运行时使用1 STT(以前的方法)和4,8 STT(我们的方法),我们使用gydF4y2BaDict_Input_SgydF4y2Bap - 1和p 2的实现。结果表明,我们的方法的运行时间(多个STT)小于前一个方法(单STT)实现p - 1和p 2。对于p - 1实现,使用4 STT比1的性能STT 18.15%, 23.41%, 100年为30.14%,5000年和50000年的模式。8 STT比1的性能STT 20%, 29.9%, 100年为35.4%,5000年和50000年的模式。图gydF4y2Ba16gydF4y2Ba显示了吞吐量性能相比之下,当我们使用前面的方法gydF4y2BaDict_Input_SgydF4y2Ba分别为p - 1和p 2实现。gydF4y2Ba

此外,数据gydF4y2Ba17gydF4y2Ba和gydF4y2Ba18gydF4y2Ba我们使用显示的吞吐量性能gydF4y2BaDict_Input_LgydF4y2Ba分别为p - 1实现和p 2实现。如两个数字所示,我们的方法优于以前的方法。事实上,性能改进我们的方法显示了类似的趋势,我们观察到在使用随机输入数据。gydF4y2Ba

5.2.2。进一步的性能优化技术的有效性gydF4y2Ba

就像前面提到的gydF4y2Ba4.3gydF4y2Ba,进一步的性能优化技术也应用在我们的实现中除了STT分区技术。gydF4y2Ba(我)gydF4y2Ba图gydF4y2Ba19gydF4y2Ba显示p 2的加速和不使用共享内存实现银行冲突。如图所示,该银行冲突影响p 2实现的性能。避免银行冲突,性能提高了1.72 ~ 4.48 x模式的数量在100 ~ 50000的范围。gydF4y2Ba(2)gydF4y2Ba为了最大化的性能优势GPU多线程功能,我们试图找到最好的/对于一个给定的数据块大小的线程数量。为此,我们进行了广泛的性能测试。p - 1的实现中,我们改变了线程的数量/同时保持相同的数据块的大小。图gydF4y2Ba20.gydF4y2Ba介绍了p - 1的运行时实现与不同数量的线程/块。结果表明,128个线程/块提供了最佳性能的所有数字的模式(100、5000和50000)。p 2实现线程的数量/块取决于共享内存的大小。因此,我们需要仔细确定共享内存的大小。(物理共享内存大小设置为48 KB在我们的实验。然而,我们设置一个逻辑共享内存大小一块线程小于48 KB考虑到多个块以多线程的方式将执行在同一硬件线程块。)通过实验,我们发现设置共享内存大小8 KB给出了最佳性能。图gydF4y2Ba21gydF4y2Ba介绍了p 2实现的结果。我们选择了线程的数量在32 ~ 512 /块。与100年和5000年模式,256个线程/块提供了最佳性能(数字gydF4y2Ba(21日)gydF4y2Ba和gydF4y2Ba21 (b)gydF4y2Ba),而512个线程/块(图50000年提供了最佳性能模式gydF4y2Ba21 (c)gydF4y2Ba)。因此,我们使用这些数字。gydF4y2Ba

5.3。性能比较XeonφgydF4y2Ba

实施Xeonφ,我们首先构建STT (s)在主机英特尔CPU。自记忆层次Xeonφ是不如GPU的内存层次结构,复杂的输入数据和STT数据都是直接复制的板载内存Xeonφ。如前所述gydF4y2Ba2gydF4y2Ba,Xeonφ两种工作模式:本机模式和卸载模式。在我们的实验中我们使用卸载模式。在卸载模式下,程序运行在主机可以选择启动或“卸载”Xeonφ协处理器的部分代码。程序员可以识别哪些行或部分的代码应该卸载,可以调用OpenMP线程。同时进行AC算法,我们将模式匹配过程Xeonφ协处理器利用协处理器的多线程功能。输入数据和STT数据从一开始,在程序执行期间没有改变。因此,他们的记忆分配和复制到协处理器只有一次在卸载阶段。此外,他们被使用在多个运行的线程之间共享gydF4y2Ba共享gydF4y2Ba条款。大量的线程创建流程模式匹配任务,每个线程处理一块输入文本。我们使用了动态调度线程之间的工作负载平衡。为了尽可能均匀分配的线程在整个系统中,应用散射亲和力。gydF4y2Ba

图gydF4y2Ba22gydF4y2Ba展示了我们的方法的加速使用4 stt Intel Xeonφ过去的方法。加速范围在1.60 ~ 2.00。加速增加而增加输入数据大小。Xeonφ结果证实的好处我们的方法来减少个人的工作集大小STT不规则的访问模式。因此,STT的分区的数量显著减少缓存缺失和导致性能改进。Xeonφ最多支持4路多线程;因此,我们可以锻炼到240个线程的实验考虑,有60计算核心。然而,当我们获得了最佳性能使用2 -或3路每核心多线程。gydF4y2Ba

5.4。STT大小和建筑成本的比较gydF4y2Ba

为了评估我们的方法的空间效率,我们测量的尺寸单大STT在前面的方法和多个小的和STT使用我们的方法生成。表gydF4y2Ba3gydF4y2Ba列出了单一STT和合并后的尺寸大小4 STT生成与不同数量的模式。100年、5000年和50000年模式,合并后的大小4 STT仅为0.88%,0.24%,和0.25%的大小大于单一STT,分别。因此,我们的方法是空间。gydF4y2Ba

我们还测量了时间建立单一STT在前面的方法和多个STT方法。运行时比较如表所示gydF4y2Ba4gydF4y2Ba。事实上,STT建筑成本增加随着STT数量的减少。例如,当模式的数量是50000,建筑成本8 STT是构建单一STT 2.22倍。因此,我们的方法是更多的时间有效地构建dfa (stt)比之前的方法。gydF4y2Ba

6。先前的研究gydF4y2Ba

交流模式匹配算法被应用于各种应用领域。事实上,网络和计算机安全,生物信息学是两个主要的AC算法广泛应用的领域。gydF4y2Ba

在网络入侵检测领域,杨和PrasannagydF4y2Ba18gydF4y2Ba)提出了一个head-body有限自动机(HBFA)方法实现了字符串模式匹配基于AC算法。HBFA实现匹配字典使用Head-DFA预定义的前缀长度。这减少了运行时内存> 20 x和性能扩展到27 x 32-core英特尔许多核心芯片。乔治Vasiliadis et al。gydF4y2Ba4gydF4y2Ba)提出了一个基于Snort入侵检测系统开源NIDS称为Gnort。GPU并行的模式匹配,他们依赖于输入数据分区间并行的线程块AC字符串匹配,而不是分割的字符串模式作为我们的方法。他们不使用共享内存中加载输入数据。相反,他们回答说在自动缓存L1缓存的GPU。史密斯et al。gydF4y2Ba19gydF4y2Ba)实现一个正则表达式匹配算法在GPU基于确定性有限自动机(扩展)。雅各Brodley [gydF4y2Ba20.gydF4y2Ba)也提出了一个解决方案将GPU签名匹配计算。他们使用了Knuth-Morris-Pratt(公里)单一字符串匹配算法代替AC算法。gydF4y2Ba

在生物信息学领域,Tumeo和别墅gydF4y2Ba8gydF4y2Ba]介绍了AC算法的高效实现对异构集群GPU加速DNA分析。咋和萨尼gydF4y2Ba5gydF4y2Ba)提出了一种并行AC算法在GPU。就像在gydF4y2Ba4gydF4y2Ba],他们瓜分了输入数据的线程块并行字符串匹配。他们使用共享内存加载输入数据;然而,他们没有考虑避免银行共享内存冲突。gydF4y2Ba

其他先前的研究多管匹配应用程序移植到IBM细胞宽带引擎(是)。Scarpazza et al。gydF4y2Ba21gydF4y2Ba,gydF4y2Ba22gydF4y2Ba细胞)移植AC-opt版本。咋et al。gydF4y2Ba23gydF4y2Ba)提出了一种压缩技术交流自动机用于在细胞。压缩AC自动机,然而,导致间接访问,得到下一个状态对于一个给定的状态和性格影响性能。别墅等。gydF4y2Ba24gydF4y2Ba)提出了一个软件AC算法的并行实现基于128 -克雷XMT处理器多线程共享内存。他们使用的特定特性XMT多线程体系结构和算法策略来最小化内存引用的数量和减少内存争用为了存档较高的性能和可伸缩性。他们还扩展这项工作通过描述AC算法的性能在各种共享内存和分布式内存架构(gydF4y2Ba6gydF4y2Ba]。gydF4y2Ba

7所示。结论gydF4y2Ba

在本文中,我们提出了一个高性能AC算法的并行化大大提高了缓存位置上不定期的DFA访问许多核心加速器芯片包括Nvidia GPU和Intel Xeonφ。我们的并行化方法分区组给定的字符串模式来生成多个设置少量的模式。然后,我们构造多个小型DFA,而不是构建单一大型DFA空间。使用多个小dfa,密集模式匹配操作同时进行对整个输入的文本字符串。这会显著减少缓存的大小的脚印STT每个核心数据的缓存,从而导致性能显著提高缓存。Nvidia Tesla甘蓝型GPU上实验结果表明我们的方法提供高达2.73倍加速与以前的方法相比使用单个大DFA和692.7 Gbps吞吐量性能。与单一CPU性能相比,我们获得了加速在127 ~ 311的范围。加速超过6 OpenMP线程并行版本上运行6个CPU核在86 ~ 183的范围。实验结果在Intel Xeonφ61 x - 86核也显示2.00倍加速比先前的方法。gydF4y2Ba

利益冲突gydF4y2Ba

作者宣称没有利益冲突有关的出版。gydF4y2Ba

确认gydF4y2Ba

作者想扩展他们的感谢匿名评论者的有价值的反馈对他们的论文。他们的评论作者大大帮助改善他们的论文的质量和演示。这项研究受到了基础科学研究项目通过韩国国家研究基金会(NRF)由教育部、科学和技术(批准号2012 r1a1a2042267)。这项工作是支持的信息与通信技术研究所由韩国政府推广(IITP)拨款(MSIP)(没有。b0101 - 15 - 0104,超级计算系统的开发进行基因组分析)。gydF4y2Ba