字符串查找是信息安全、信息过滤领域的重要操作，尤其是对大文本的实时处理。这篇作为实例，使用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)：
1 int compare(__global const uchar* text, __local const uchar* pattern, uint length){ 2 for(uint l=0; l lastSearchIdx) 29 return;30 if(endSearchIdx > lastSearchIdx) 31 endSearchIdx = lastSearchIdx;32 33 for(int idx = localIdx; idx < patternLength; idx+=localSize)34 localPattern[idx] = pattern[idx];35 barrier(CLK_LOCAL_MEM_FENCE);36 37 for(uint stringPos=beginSearchIdx+localIdx; stringPos 3.2.tool.h 、tool.cpp
见：
3.3 StringSearch.cpp
1 #include2 #include "tool.h" 3 #include 4 #include 5 #include 6 #include 7 #include 8 #include 9 using namespace std; 10 11 12 int main(int argc, char* argv[]) 13 { 14 cl_int status; 15 /**Step 1: Getting platforms and choose an available one(first).*/ 16 cl_platform_id platform; 17 getPlatform(platform); 18 19 /**Step 2:Query the platform and choose the first GPU device if has one.*/ 20 cl_device_id *devices=getCl_device_id(platform); 21 22 /**Step 3: Create context.*/ 23 cl_context context = clCreateContext(NULL,1, devices,NULL,NULL,NULL); 24 25 /**Step 4: Creating command queue associate with the context.*/ 26 cl_command_queue commandQueue = clCreateCommandQueue(context, devices[0], 0, NULL); 27 28 /**Step 5: Create program object */ 29 const char *filename = "string_search_kernel.cl"; 30 string sourceStr; 31 status = convertToString(filename, sourceStr); 32 const char *source = sourceStr.c_str(); 33 size_t sourceSize[] = {strlen(source)}; 34 cl_program program = clCreateProgramWithSource(context, 1, &source, sourceSize, NULL); 35 36 /**Step 6: Build program. */ 37 status=clBuildProgram(program, 1,devices,NULL,NULL,NULL); 38 39 40 /**Step 7: Initial input,output for the host and create memory objects for the kernel*/ 41 string textStr; //StringSearch_Input.txt 42 convertToString("StringSearch_Input.txt", textStr); 43 const char * text = textStr.c_str(); 44 int textlen=strlen(text); 45 46 char * pattern="info"; 47 int patternlen=strlen(pattern); 48 int maxSearchLength=256*64; 49 int * resultCount=new int[1]; 50 *resultCount=0; 51 int * result=new int[textlen]; 52 memset(result,0,sizeof(int)*textlen); 53 54 cl_mem textBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, sizeof(char)*textlen,(void *)text, NULL); //global memory 55 cl_mem patternBuffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY|CL_MEM_COPY_HOST_PTR ,sizeof(char)*patternlen, (void *)pattern, NULL); 56 cl_mem resultCountBuffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY|CL_MEM_COPY_HOST_PTR ,sizeof(int), (void *)resultCount, NULL); 57 cl_mem resultBuffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY|CL_MEM_COPY_HOST_PTR ,sizeof(int)*textlen, (void *)result, NULL); 58 59 /**Step 8: Create kernel object */ 60 cl_kernel kernel = clCreateKernel(program,"StringSearch", NULL); 61 62 /**Step 9: Sets Kernel arguments.*/ 63 status = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&textBuffer); //global 64 status = clSetKernelArg(kernel, 1, sizeof(int), &textlen); //private 65 status = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&patternBuffer); //global 66 status = clSetKernelArg(kernel, 3, sizeof(int), &patternlen); //private 67 status = clSetKernelArg(kernel, 4, sizeof(int), &maxSearchLength); //private 68 status = clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *)&resultCountBuffer); //global 69 status = clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *)&resultBuffer); //global 70 status = clSetKernelArg(kernel, 7, sizeof(char)*patternlen, NULL); //local 71 72 /**Step 10: Running the kernel.*/ 73 cl_event enentPoint; 74 int globalWorkItem=textlen/64; 75 76 if(textlen%64 != 0) 77 globalWorkItem++; 78 size_t groupNUm[1]={globalWorkItem}; 79 size_t localNUm[1]={ 256}; 80 81 status = clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, groupNUm, localNUm, 0, NULL, &enentPoint); 82 83 clWaitForEvents(1,&enentPoint); ///wait 84 clReleaseEvent(enentPoint); 85 int count=0; 86 status = clEnqueueReadBuffer(commandQueue, resultCountBuffer, CL_TRUE, 0, sizeof(int), &count, 0, NULL, NULL); 87 cout<<"\nNumber of matches:"< <
本文：