基于GPU加速的快速字符串匹配算法

续士强　祝永志

摘要：如何提高字符串实际匹配效率一直是信息匹配领域中非常重要的研究课题。在分析字符串匹配并行规律的基础上，结合GPU并行体系结构，对Sunday算法实现并行化。在CPU和GPU不同计算平台上分别做了对比实验，实验结果显示基于GPU并行实现的Sunday算法比传统Sunday算法具有更高的匹配速度。

关键词关键词：CUDA；Sunday算法；并行计算

DOIDOI：10.11907/rjdk.143766

中图分类号：TP312

文献标识码：A文章编号文章编号：16727800（2015）002005103

基金项目基金项目：山东省自然科学基金项目（ZR2013FL015）；山东省研究生教育创新资助计划项目（SDYY12060）

作者简介作者简介：续士强（1990-），男，山东滕州人，曲阜师范大学信息科学与工程学院硕士研究生，研究方向为网络与并行计算；祝永志（1964-），男，吉林榆树人，曲阜师范大学信息科学与工程学院教授、硕士生导师，研究方向为网络与并行计算。

0引言

字符串匹配被广泛应用于解决搜索引擎、网络入侵检测、计算机病毒特征码匹配以及DNA序列匹配等诸多问题中。目前关于字符串匹配算法的研究已经相对成熟，传统的串匹配算法实现思想主要分为前缀搜索、后缀搜索和子串搜索，经典的算法[1]有KMP算法、Shift-And算法、BM算法、BDM算法等等。然而传统算法在实际应用中并没有表现理论上的优势，比如，在一个普通文本文件中查找字符串，BF算法用的时间通常比KMP算法要少，即使后者在时间复杂度上要优于前者。因此，选择一种在实际应用中性能较好的算法，对于提高效率至关重要。

随着图形处理器性能的不断提高，特别是通用并行计算架构（CUDA）的推出，图形处理器已不再局限于图形图像处理，其在通用计算方面已经远远超越了CPU。由于GPU的特点就是处理密集型数据和并行数据计算，因此，CUDA编程模型非常适合解决大规模并行计算问题。

1CUDA编程模型

CUDA（Compute Unified Device Architecture），是由NVIDIA推出的通用并行计算架构，它在C语言基础上添加了适用于GPU并行计算的API和开发库，是在GPU强大计算能力基础上建立的效率更高的密集数据计算解决方案。

在NVIDIA的GPU中，最基本的处理单元是SP（流处理器），多个SP组成一个SM（流多处理器），它们之间可以共享控制逻辑和指令缓存。在CUDA编程下，程序执行的最小单位是thread（线程），每个thread都有各自的一份register和local memory。一组thread组成一个block（线程块），在同一个block中的thread可以存取同一块共享内存，最后执行同一kernel函数的block构成一个grid（块网格）。CUDA架构下的程序包括运行在CPU上负责串行执行的Host端，以及运行在GPU上作并行计算的Device端。Kernel函数由host端调用，能生成大量的线程，每个线程处理对应的数据块，从而实现了数据的并行处理。grid的所有线程都执行同一个kernel函数。CUDA的device在实际执行时，以block为单位，并将每个block分给SM执行，同一block中的thread，又以32个组成一个warp一起执行。每个SM只执行一个block中的warp，由于在实际执行过程中，某个warp需要访问全局存储器，这时SM就会切换到别的warp来继续执行。因此，一个SM可以同时处理多个block。

2基于GPU的Sunday算法并行化

2.1Sunday字符串匹配算法

Sunday算法源于BM算法充分利用坏字符的思想。Sunday算法将计算好后缀的时间节约下来，将所有情况归为坏字符规则来匹配，这对短模式串匹配十分有利。因此，Sunday算法的匹配过程基本应用的就是BM算法的坏字符规则，即当字符串不匹配时，判断目标串进行匹配的下一字符是否在模式串中出现，如果出现，则模式串移动到目标串中最右端的该字符串的下一位置；如果没有出现，则模式串直接跳过模式串字符长度。

Sunday算法实现如下：

while（main_i

{

int tem=main_i；

while（sub_j

{

if（mainStr[main_i] == subStr[sub_j]）

{

main_i++；

sub_j++；

continue；

}

else{

if（tem+subStrLen > mainStrLen）

return -1；

char firstRightChar=

mainStr[tem+subStrLen]；

main_i+=

char step[（unsigned char）firstRightChar]；

sub_j=0；

break；

}

}

if（sub_j == subStrLen）

return main_i-subStrLen；

}

2.2字符串并行划分设计

CUDA编程模型是基于并行执行多个线程来实现大规模快速计算的，因此字符串要实现并行匹配，就必须将原有的字符串划分成多个小字符串，再将小字符串分散到各个线程中，线程对分配的小字符串进行匹配。但是，仅仅做到简单划分是不够的，通常串行的字符串匹配算法，无论算法如何设计向前跳跃，匹配始终带有一定的顺序性，在划分后的局部字符串匹配位置必须到下一小字符串才能找到。为了解决划分字符串过程中存在的漏解问题，采用冗余数据是简单可行的方法。

假设将长度为T的文本串均匀分成互不重叠的m段，分布在0～（m-1）线程中，则显然每个线程中字符串长度为[n/m]，再将长度为s的模式串分配到各个线程，使得每个线程都对分到的小字符串进行匹配[3]。为了解决漏解问题，在每个线程分得的小字符串末尾位置存储与其相邻的长度为s-1的字符串。

2.3Sunday算法并行实现

串行实现的Sunday算法，通过极少的比较次数来获得较大的跳跃距离，与BF算法相比，避免了大量回溯匹配的产生。而基于GPU这种并行体系结构的并行算法，利用多个线程并行执行解决问题，原有较长的字符串划分到各个线程，使得每个线程字符串长度明显缩短，比串行集中式的匹配算法速度提高很多。

2.3.1数据预处理

为了最大限度地利用CPU资源，目标串的存取与分片工作放在CPU进行。首先在内存中分配一块大致与目标串物理存储大小相等的存取空间。比如目标串存储的文本大小为10M，那么就在内存中申请大约10M的地址空间，目的是为了使目标串读入主存，便于程序存取。当然，实际申请地址空间时，空间大小是根据读取到的目标串大小而设定的。通过I/O操作将目标串从外部文件读入到内存后，对目标串进行如下分片操作：

（1）根据目标串的总长度与CUDA的线程数来确定每个线程要处理字符串的长度。虽然GPU可以创建上百万甚至上千万线程，但是同时并行的线程总数却十分有限。因此，创建合理数量的线程不仅可以避免一些额外开销，而且可以提升算法总体效率。以GT750M为例，每个block最多可以同时运行的线程数是2 048，但根据实验结果，取1 024比较好。这是因为每个CUDA线程处理性能有限，分片后小字符串如果太长，线程匹配耗时增加，算法的整体匹配效率会变差。如果太短，那么分配的线程总数就多，用于线程分配和回收的额外时间大大增加，算法的效率就得不到提高。因此，合理的分片对于最终匹配效果十分重要。经过实验，分片后的小字符串长度验证取64为宜。

（2）字符串的漏解处理[4]。由于对字符串进行了分片操作，小字符串末尾和开始位置的字符就丧失了原来的顺序，同一串的字符分到不同线程，而这一字符串可能恰好是匹配串，这样就造成匹配过程的漏解情况。为了解决这个问题，必须在每个小字符串后面添加下一段字符串的前m-1个的字符（m为模式串长度）。这里需要说明的是，添加冗余数据的处理在CPU端进行。由于模式串的长度相对较短，即使在分片较多的情况下，所添加的冗余字符对算法的匹配效率也不会造成大的影响。需要注意的是，添加冗余字符后，模式串的实际匹配位置需要作移位处理，保持结果正确。

2.3.2并行匹配

改进算法对字符串实际匹配操作都是在GPU中完成的，因此目标串要再从内存读取到显存，与内存申请地址空间类似。通过cudaMalloc函数在显存中申请地址空间，然后再调用cudaMemcpy函数将内存中的目标串复制到显存中。这样，GPU就能直接对显存中的目标串进行匹配操作。调用kernel函数完成字符的匹配后，再把显存存储的匹配结果复制到内存的变量中进行显示。

Kernel函数对字符串匹配操作过程：①根据blockIdx与threadIdx变量确定当前线程的实际线程号；②根据线程号找到要处理的字符串位置；③对每个小字符串实现Sunday算法。每个线程处理字符串长度与分片后的小字符串长度一致。

Kernel函数的部分代码如下：

__global__ void sunday_kernel（char *g_subStr，int subStr_len，int aim_len，char *aim_text，int *gpu_pos）

{

int i=threadIdx.x+threadIdx.y*16+blockIdx.x*256；

int j=0；

int c = 0； //计数

int m = i*64；

int n = 0；

while（m+n

if（aim_text[m+n]== g_subStr [j]）

{

n++；

j++；

count++；

}

else {

n++；

j = 0；

count = 0；

}

if（c== subStr_len）

{

*gpu_pos = m+n-count；

}

}

3实验结果与分析

实验环境硬件配置：CPU：Intel Core i5-3230M，主频：2.60GHz，内存：4GB。GPU采用GT750M 核心频率950MHz，显存2G。操作系统是Windows XP professional 2002 Service Package 3，软件平台为CUDA 5.5，实验集成环境Visual Studio 2010。

使用人类基因DNA部分序列集合字符串作为文本串匹配数据集。本文第一组实验，测试在相同文本大小的目标串情况下，不同匹配算法消耗的匹配时间对比，实验文本大小为30M，模式串为“test”，实验结果见表1，匹配时间结果以ms为单位；第二组实验，测试CPU与GPU固定模式串长度在不同文本大小字符串的匹配时间对比，实验过程中对字符串文本随机插入模式串“test”，再对模式串进行匹配，实验结果见表2，匹配时间以ms为单位，匹配位置单位为字符个数。

表1不同算法消耗的匹配时间对比

[]BF算法[]KMP算法[]Sunday算法

匹配时间（ms）[]85.784 66[]215.690 48[]73.988 68

匹配位置[]32 157 335[]32 157 335[]32 157 335

表2Sunday算法与GPU并行Sunday算法字符串匹配消耗时间对比

文本大小[]1M[]8M[]16M[]30M[]53M[]77M

CPU匹配时间（ms）[]2.942 78[]20.341 69[]40.153 63[]75.152 00[]133.682 68[]191.519 03

GPU匹配时间（ms）[]2.008 25[]7.455 17[]14.277 82[]25.333 25[]43.330 85[]61.758 17

匹配匹配位置[]1 024 997[]8 353 018[]17 037 000[]32 157 327[]55 622 541[]80 752 081

从第一组实验结果可以看出，在相同文本大小的目标串和相同的模式串情况下，Sunday算法的匹配时间最短，BF算法次之，而KMP算法匹配时间最长。传统的字符串匹配算法实现的基本思路是，根据模式串按照从左到右或从右到左的顺序依次与模式串匹配，若不匹配，则根据相应规则进行移位跳跃。大部分的字符串匹配算法只是对如何进行移位跳跃进行改进，而这个改进还必须根据模式串的自身结构是否具有前后缀匹配来实现。实际应用中，特别是文本字符串匹配时，输入的模式串很难经常有前后缀自相匹配的情况。因此，针对于这方面的改进算法，匹配时间比BF算法没有明显的降低，甚至比BF算法耗费的时间更长，这是因为改进算法在移动跳跃之前，还要对模式串遍历和比较。而Sunday算法通过模式串尾部字符与目标串相应字符比较，获得尽可能大的跳跃距离，避免了重复字符匹配。因此，Sunday算法大大提高了匹配效率。

由表2实验结果可以看出，在不同大小的文本字符串匹配中，GPU并行的Sunday算法与传统Sunday算法维持约3倍的加速比。GPU先天强大的并行处理能力为字符串并行匹配提供了良好的工具。在实际程序运行中，GPU线程的分配和结果数据的回传占用了较多的时间，尽管如此，字符串的匹配速度还是获得了较大的提升。与此同时，将模式串存储到共享内存的方式也避免了频繁访问显存带来的额外时间开销，因为显存的读写速度比共享内存的速度慢得多。

4结语

本文在分析字符串匹配并行执行特点的基础上，利用GPU数据并行化的架构特点，设计出适合于GPU运行的快速字符串匹配算法。虽然算法使得字符串匹配速度有所提升，但是还有较大的改进空间，将字符串匹配放在显卡上进行还可以进一步提高字符读取速度，提高字符匹配效率。

参考文献参考文献：

＼[1＼]谷岳，谷建华.基于GPU加速的并行字符串匹配算法[J].微电子学与计算机，2013（9）：3033.

[2]潘冠华，张兴忠.Sunday算法效率分析[J].计算机应用，2012（11）：30823084，3088.

[3]陈国良，林洁，顾乃杰. 分布式存储的并行串匹配算法的设计与分析[J].软件学报，2000（6）： 771778.

[4]唐定车，刘任任，谭建龙.基于统一计算设备架构的并行串匹配算法[J].计算机应用，2009（1）： 399401.

[5][美]科克，胡文美.大规模并行处理器编程实战[M].北京：清华大学出版社，2010：3177.

责任编辑（责任编辑：杜能钢）