字符串查找是信息安全、信息过滤领域的重要操作，尤其是对大文本的实时处理。这篇作为实例，使用GPU OpenCL进行精确模式串查找。
1.加速方法
(1)将少量常量数据，如模式串长度、文本长度等，保存在线程的private memory中。
(2)将模式串保存在GPU的local memory中，加速线程对模式串的访问。
(3)将待查找的文本保存在global memory中，使用尽可能多线程访问global memory，减小线程平均访存时间。
(4)每个work-group中的线程操作文本中一段，多个work-group并行处理大文本。
2.同步
(1)work-group内，使用CLK_LOCAL_MEM_FENCE、CLK_GLOBAL_MEM_FENCE
(2)全局使用对__global int 的原子操作，来保证每个线程将结果写到全局内存的正确位置。设备支持的操作可以通过查询设备的扩展获得，如下图，可知核函数支持原子操作、printf操作：
3.代码实例，大文本精确模式串搜索
3.1 核函数(string_search_kernel.cl)：
int compare(__global const uchar* text, __local const uchar* pattern, uint length){
for(uint l=; l<length; ++l){
if (text[l] != pattern[l])
return ;
}
return ;
}
__kernel void
StringSearch (
__global uchar* text, //Input Text
const uint textLength, //Length of the text
__global const uchar* pattern, //Pattern string
const uint patternLength, //Pattern length
const uint maxSearchLength, //Maximum search positions for each work-group
__global int* resultCount, //Result counts (global)
__global int* resultBuffer, //Save the match result
__local uchar* localPattern) //local buffer for the search pattern
{
int localIdx = get_local_id();
int localSize = get_local_size();
int groupIdx = get_group_id();
uint lastSearchIdx = textLength - patternLength + ;
uint beginSearchIdx = groupIdx * maxSearchLength;
uint endSearchIdx = beginSearchIdx + maxSearchLength;
if(beginSearchIdx > lastSearchIdx)
return;
if(endSearchIdx > lastSearchIdx)
endSearchIdx = lastSearchIdx;
for(int idx = localIdx; idx < patternLength; idx+=localSize)
localPattern[idx] = pattern[idx];
barrier(CLK_LOCAL_MEM_FENCE);
for(uint stringPos=beginSearchIdx+localIdx; stringPos<endSearchIdx; stringPos+=localSize){
if (compare(text+stringPos, localPattern, patternLength) == ){
int count = atomic_inc(resultCount);
resultBuffer[count] = stringPos;
//printf("%d ",stringPos);
}
barrier(CLK_LOCAL_MEM_FENCE);
}
}3.2.tool.h 、tool.cpp
见：http://www.cnblogs.com/xudong-bupt/p/3582780.html
3.3 StringSearch.cpp
#include <CL/cl.h>
#include "tool.h"
#include <string.h>
#include <stdio.h>
#include <stdlib.h>
#include <iostream>
#include <string>
#include <fstream>
using namespace std; int main(int argc, char* argv[])
{
cl_int status;
/**Step 1: Getting platforms and choose an available one(first).*/
cl_platform_id platform;
getPlatform(platform); /**Step 2:Query the platform and choose the first GPU device if has one.*/
cl_device_id *devices=getCl_device_id(platform); /**Step 3: Create context.*/
cl_context context = clCreateContext(NULL,, devices,NULL,NULL,NULL); /**Step 4: Creating command queue associate with the context.*/
cl_command_queue commandQueue = clCreateCommandQueue(context, devices[], , NULL); /**Step 5: Create program object */
const char *filename = "string_search_kernel.cl";
string sourceStr;
status = convertToString(filename, sourceStr);
const char *source = sourceStr.c_str();
size_t sourceSize[] = {strlen(source)};
cl_program program = clCreateProgramWithSource(context, , &source, sourceSize, NULL); /**Step 6: Build program. */
status=clBuildProgram(program, ,devices,NULL,NULL,NULL); /**Step 7: Initial input,output for the host and create memory objects for the kernel*/
string textStr; //StringSearch_Input.txt
convertToString("StringSearch_Input.txt", textStr);
const char * text = textStr.c_str();
int textlen=strlen(text); char * pattern="info";
int patternlen=strlen(pattern);
int maxSearchLength=*;
int * resultCount=new int[];
*resultCount=;
int * result=new int[textlen];
memset(result,,sizeof(int)*textlen); cl_mem textBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, sizeof(char)*textlen,(void *)text, NULL); //global memory
cl_mem patternBuffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY|CL_MEM_COPY_HOST_PTR ,sizeof(char)*patternlen, (void *)pattern, NULL);
cl_mem resultCountBuffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY|CL_MEM_COPY_HOST_PTR ,sizeof(int), (void *)resultCount, NULL);
cl_mem resultBuffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY|CL_MEM_COPY_HOST_PTR ,sizeof(int)*textlen, (void *)result, NULL); /**Step 8: Create kernel object */
cl_kernel kernel = clCreateKernel(program,"StringSearch", NULL); /**Step 9: Sets Kernel arguments.*/
status = clSetKernelArg(kernel, , sizeof(cl_mem), (void *)&textBuffer); //global
status = clSetKernelArg(kernel, , sizeof(int), &textlen); //private
status = clSetKernelArg(kernel, , sizeof(cl_mem), (void *)&patternBuffer); //global
status = clSetKernelArg(kernel, , sizeof(int), &patternlen); //private
status = clSetKernelArg(kernel, , sizeof(int), &maxSearchLength); //private
status = clSetKernelArg(kernel, , sizeof(cl_mem), (void *)&resultCountBuffer); //global
status = clSetKernelArg(kernel, , sizeof(cl_mem), (void *)&resultBuffer); //global
status = clSetKernelArg(kernel, , sizeof(char)*patternlen, NULL); //local /**Step 10: Running the kernel.*/
cl_event enentPoint;
int globalWorkItem=textlen/; if(textlen% != )
globalWorkItem++;
size_t groupNUm[]={globalWorkItem};
size_t localNUm[]={}; status = clEnqueueNDRangeKernel(commandQueue, kernel, , NULL, groupNUm, localNUm, , NULL, &enentPoint); clWaitForEvents(,&enentPoint); ///wait
clReleaseEvent(enentPoint);
int count=;
status = clEnqueueReadBuffer(commandQueue, resultCountBuffer, CL_TRUE, , sizeof(int), &count, , NULL, NULL);
cout<<"\nNumber of matches:"<<count<<endl; /**Step 12: Clean the resources.*/
status = clReleaseKernel(kernel);//*Release kernel.
status = clReleaseProgram(program); //Release the program object.
status = clReleaseMemObject(resultBuffer);//Release mem object.
status = clReleaseMemObject(textBuffer);//Release mem object.
status = clReleaseMemObject(resultCountBuffer);//Release mem object.
status = clReleaseMemObject(patternBuffer);//Release mem object.
status = clReleaseCommandQueue(commandQueue);//Release Command queue.
status = clReleaseContext(context);//Release context. free(devices);
free(result);
free(resultCount); getchar();
return ;
}