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

李　磊

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

[關鍵詞]GPU BF CUDA

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

一、引言

NVIDIA公司在1999年發布GeForce 256圖形處理芯片時首先提出GPU（Graphic Processing 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<=n-m；i++)

{ j=0； k=i ； //從第i個位置開始搜索字符串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<=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()； //同步函數，所有線程到這步等待同步

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．