论文部分内容阅读
[摘要]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．
[关键词]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．