利用ＧＰＵ进行字符串匹配
作者：李　磊
来源：《硅谷》2009年第02期
        [摘要]GPU通过SIMD（Single Instruction Multiple Data，单指令多数据）对图像数据进行并行处理。字符串的匹配在信息检索、计算机病毒码匹配和生物基因技术领域中都有应用。探讨利用GPU进行字符串的匹配。
        [关键词]GPU BF CUDA
        中图分类号：TP3文献标识码：A文章编号：1671－7597（2009）0120055－01
       
        一、引言
       
        NVIDIA公司在1999年发布GeForce 256图形处理芯片时首先提出GPU（Graphic Process
ing Unit，图形处理芯片）的概念。GPU的32位的浮点渲染精度、向量处理特征和超长流水线结构特点，使它具有对密集性数据的计算能力，在通用计算机上提供一种并行平台。目前，GPU在分布式计算、生物制药、天气预报等非图形数据处理领域都有广泛应用。
        字符串匹配在信息检索、计算机病毒码匹配和生物基因技术领域中都有应用。字符串的匹配算法有很多，比如BF算法、KMP算法、RK算法。
        CUDA（Compute Unified Device Architecture，统一计算设备架构）是用于GPU计算的开发环境，它是一个全新的软硬件架构，可以将GPU视为一个并行数据计算的设备，对所进行的计算进行分配和管理。CUDA的GPU编程语言基于标准的C语言，对大多数程序员来说还是很容易掌握的。
        本文就BF算法在GPU环境下利用CUDA软件开发工具的实现展开讨论。
       
        二、介绍一下CUDA编程的特点
       
        通过CUDA编程时，将GPU看作可以并行执行多个线程的计算设备(compute device)。它作为主CPU的协处理器来运作。应用程序首先被主CPU执行过程，而应用程序数据并行的、计算密集的部分给GPU执行，GPU最后把计算的结果返回给主CPU中的应用程序。即GPU执行的部分函数化。要达到这种效果，可以将这样一个函数编译到GPU设备的指令集合中，并将得到的程序（叫做内核，kernel）下载到GPU设备上。CPU和GPU都保留自己的DRAM，分别称为主机内存(host memory)和设备内存(device memory)。用户可以通过优化的API调用将数据从一个DRAM复制到其他DRAM中，而优化的API调用使用了设备的高性能直接内存访问(DMA)引擎。
        下面是在主CPU下完成的BF朴素匹配算法，C语言描述如下：
        int BFString(char* query , char* subject,)
        { int i , j , k , num= -1；
        int m=strlen(query)； //模式串长度
        int n=strlen(subject)； //目标串长度
        for(i=0； i
        { j=0； k=i ； //从第i个位置开始搜索字符串字符串长度17模式串长度subject，看是否
        存在子字符串跟模式串query一样
        while( j< m && subject[k]= = query[j]){k++； j++； }
        if(j= =m) num++；//条件成立，则到模式串，记录个数加1
        }
        return num； //返回个数
        }
        这个函数里面的形式参数、局部变量都是在内存中定义的，CPU可以直接访问。
        为了体现GPU的并行性，我们把字符串的每次比较交给不同的线程完成。比如thread_0从i=0位置开始，thread_1从i=1位置开始，thread_2从i=2位置开始。。。。。。thread_n从i
=n位置开始（假设用n+1个线程来执行）；然后thread_0从i=n+1位置开始，thread_1从i=n+2位置开始。。。。这样循环直到所有的搜索都完成。目前GeForce 8800GT一个块(block)最多只能有512个threads。假设一个kernel只由一个block组成。这里需要在CUDA的初始化文件，加入下面的#define语句，定义thread数目。
        #define THREAD_NUM 256
        接着把kernel程序改成如下形式：
        __global__ static void BFString(char * query , char * subject ,int number ，int len_que , int len_sub )
        { //query模式串，subject目标串
        extern __shared__ int shared[];//共享内存空间，保存匹配的个数
        const int tid = threadIdx.x； //获得线程的id
        // const int bid = blockIdx.x;获得block的id,此时block只有一块，故省略
        int i , j , k ；shared[id]=0;
        for(i=tid ；i
        { j=0；k=i ；
        while( j< len_que && subject[k]= = query[j]){k++；j++；}
        if(j= =len_que) shared[tid]++；
        }
        __syncthreads()； //同步函数，所有线程到这步等待同步
        if(tid = = 0) {
        for(i = 1； i < THREAD_NUM； i++) shared[0] += shared[tid]；
        number = shared[0]；//利用线程0来计算匹配的子字符串的个数
        }
        }
        这个函数里面的形式参数和其它变量是在GPU的存储器里面。那么在main()程序中，就需要利用CUDA的cudamalloc()分配GPU的存储器空间和cudamemcpy()复制主存的内容到显存。调用kernel函数后，还要将空间释放（用cudafree()函数）和将number的值复制回主存（用cudamemcpy()函数）。函数的具体操作请查看CUDA相应的文档。
        由于1个block中的thread的个数是有限制的，要想有更多的thread来参与计算，就必须增加block的数目。注意，一个block里面的线程有一个shared memory。block之间的shared memory不能相互访问。比如，我们这个时候在#define THREAD_NUM的位置加上：
        #define BLOCK_NUM 32
        则kernel程序中需要修改的主要是在for循环：
        for(i = bid * THREAD_NUM + tid； i < len_sublen_que ；
        i += BLOCK_NUM * THREAD_NUM)
        { j=0；k=i ；
        while( j< len_que && subject[k]= = query[j]){k++；j++；}
        if(j= =len_que) shared[tid]++；
        }
        这里要注意一点，就是复制结果回去的时候要考虑到BLOCK_NUM个shared[0]。
       
        参考文献：
        [1]张庆丹、戴正华、冯圣中、孙凝晖，基于GPU的串匹配算法研究，计算机应用[Ｊ]．2006．26(7):1735-1737．

版权声明：本站内容均来自互联网，仅供演示用，请勿用于商业和其他非法用途。如果侵犯了您的权益请与我们联系QQ:729038198，我们将在24小时内删除。