利用ＧＰＵ進行字符串匹配

李　磊

[摘要]GPU通過SIMD（Single Instruction Multiple Data，單指令多數(shù)據(jù)）對圖像數(shù)據(jù)進行并行處理。字符串的匹配在信息檢索、計算機病毒碼匹配和生物基因技術(shù)領(lǐng)域中都有應(yīng)用。探討利用GPU進行字符串的匹配。

[關(guān)鍵詞]GPU BF CUDA

中圖分類號：TP3文獻標(biāo)識碼：A文章編號：1671－7597（2009）0120055－01

一、引言

NVIDIA公司在1999年發(fā)布GeForce 256圖形處理芯片時首先提出GPU（Graphic Processing Unit，圖形處理芯片）的概念。GPU的32位的浮點渲染精度、向量處理特征和超長流水線結(jié)構(gòu)特點，使它具有對密集性數(shù)據(jù)的計算能力，在通用計算機上提供一種并行平臺。目前，GPU在分布式計算、生物制藥、天氣預(yù)報等非圖形數(shù)據(jù)處理領(lǐng)域都有廣泛應(yīng)用。

字符串匹配在信息檢索、計算機病毒碼匹配和生物基因技術(shù)領(lǐng)域中都有應(yīng)用。字符串的匹配算法有很多，比如BF算法、KMP算法、RK算法。

CUDA（Compute Unified Device Architecture，統(tǒng)一計算設(shè)備架構(gòu)）是用于GPU計算的開發(fā)環(huán)境，它是一個全新的軟硬件架構(gòu)，可以將GPU視為一個并行數(shù)據(jù)計算的設(shè)備，對所進行的計算進行分配和管理。CUDA的GPU編程語言基于標(biāo)準(zhǔn)的C語言，對大多數(shù)程序員來說還是很容易掌握的。

本文就BF算法在GPU環(huán)境下利用CUDA軟件開發(fā)工具的實現(xiàn)展開討論。

二、介紹一下CUDA編程的特點

通過CUDA編程時，將GPU看作可以并行執(zhí)行多個線程的計算設(shè)備(compute device)。它作為主CPU的協(xié)處理器來運作。應(yīng)用程序首先被主CPU執(zhí)行過程，而應(yīng)用程序數(shù)據(jù)并行的、計算密集的部分給GPU執(zhí)行，GPU最后把計算的結(jié)果返回給主CPU中的應(yīng)用程序。即GPU執(zhí)行的部分函數(shù)化。要達到這種效果，可以將這樣一個函數(shù)編譯到GPU設(shè)備的指令集合中，并將得到的程序（叫做內(nèi)核，kernel）下載到GPU設(shè)備上。CPU和GPU都保留自己的DRAM，分別稱為主機內(nèi)存(host memory)和設(shè)備內(nèi)存(device memory)。用戶可以通過優(yōu)化的API調(diào)用將數(shù)據(jù)從一個DRAM復(fù)制到其他DRAM中，而優(yōu)化的API調(diào)用使用了設(shè)備的高性能直接內(nèi)存訪問(DMA)引擎。

下面是在主CPU下完成的BF樸素匹配算法，C語言描述如下：

int BFString(char* query , char* subject,)

{ int i , j , k , num= -1；

int m=strlen(query)； //模式串長度

int n=strlen(subject)； //目標(biāo)串長度

for(i=0； i<=n-m；i++)

{ j=0； k=i ； //從第i個位置開始搜索字符串subject，看是否

存在子字符串跟模式串query一樣

while( j< m && subject[k]= = query[j]){k++； j++； }

if(j= =m) num++；//條件成立，則找到模式串，記錄個數(shù)加1

}

return num； //返回個數(shù)

}

這個函數(shù)里面的形式參數(shù)、局部變量都是在內(nèi)存中定義的，CPU可以直接訪問。

為了體現(xiàn)GPU的并行性，我們把字符串的每次比較交給不同的線程完成。比如thread_0從i=0位置開始，thread_1從i=1位置開始，thread_2從i=2位置開始。。。。。。thread_n從i=n位置開始（假設(shè)用n+1個線程來執(zhí)行）；然后thread_0從i=n+1位置開始，thread_1從i=n+2位置開始。。。。這樣循環(huán)直到所有的搜索都完成。目前GeForce 8800GT一個塊(block)最多只能有512個threads。假設(shè)一個kernel只由一個block組成。這里需要在CUDA的初始化文件，加入下面的#define語句，定義thread數(shù)目。

#define THREAD_NUM 256

接著把kernel程序改成如下形式：

__global__ static void BFString(char * query , char * subject ,int number ，int len_que , int len_sub )

{ //query模式串，subject目標(biāo)串

extern __shared__ int shared[];//共享內(nèi)存空間，保存匹配的個數(shù)

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<=len_sublen_que ； i +=THREAD_NUM)

{ j=0；k=i ；

while( j< len_que && subject[k]= = query[j]){k++；j++；}

if(j= =len_que) shared[tid]++；

}

__syncthreads()； //同步函數(shù)，所有線程到這步等待同步

if(tid = = 0) {

for(i = 1； i < THREAD_NUM； i++) shared[0] += shared[tid]；

number = shared[0]；//利用線程0來計算匹配的子字符串的個數(shù)

}

}

這個函數(shù)里面的形式參數(shù)和其它變量是在GPU的存儲器里面。那么在main()程序中，就需要利用CUDA的cudamalloc()分配GPU的存儲器空間和cudamemcpy()復(fù)制主存的內(nèi)容到顯存。調(diào)用kernel函數(shù)后，還要將空間釋放（用cudafree()函數(shù)）和將number的值復(fù)制回主存（用cudamemcpy()函數(shù)）。函數(shù)的具體操作請查看CUDA相應(yīng)的文檔。

由于1個block中的thread的個數(shù)是有限制的，要想有更多的thread來參與計算，就必須增加block的數(shù)目。注意，一個block里面的線程有一個shared memory。block之間的shared memory不能相互訪問。比如，我們這個時候在#define THREAD_NUM的位置加上：

#define BLOCK_NUM 32

則kernel程序中需要修改的主要是在for循環(huán)：

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]++；

}

這里要注意一點，就是復(fù)制結(jié)果回去的時候要考慮到BLOCK_NUM個shared[0]。

參考文獻：

[1]張慶丹、戴正華、馮圣中、孫凝暉，基于GPU的串匹配算法研究，計算機應(yīng)用[Ｊ]．2006．26(7):1735-1737．