GPGPU OpenCL How to achieve accurate string lookup?

1. Acceleration Method

(1) Save a small number of constant data, such as pattern string length, text length, and so on, in the private memory of the thread.

(2) The mode string is saved in the local memory of the GPU, which accelerates the thread's access to the pattern string.

(3) Save the text to be found in the global memory, using as many threads as possible to access the global memory, reducing the average time to visit the thread.

(4) The threads in each work-group manipulate text in one paragraph, and multiple Work-group handle large text in parallel.

　　2. Sync

(1) in Work-group, using Clk_local_mem_fence, clk_global_mem_fence

(2) Global use of atomic operations on __global int to ensure that each thread writes the result to the correct location of the global memory. The operation supported by the device can be obtained by the extension of the query device, as shown in the following figure, the kernel function supports atomic operation, printf operation:

　　3. Code instance, large text exact pattern string search

3.1 Kernel function (string_search_kernel.cl):

int compare (__global const uchar* text, __local const uchar* pattern, uint length) {

for (UINT l=0 L)

if (Text[l]!= pattern[l])

return 0;

}

return 1;

}

__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 (0);

int localsize = get_local_size (0);

int groupidx = get_group_id (0);

UINT LASTSEARCHIDX = textlength-patternlength + 1;

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

if (compare (Text+stringpos, Localpattern, patternlength) = 1) {

int count = Atomic_inc (Resultcount);

Resultbuffer[count] = Stringpos;

printf ("%d", stringpos);

}

Barrier (clk_local_mem_fence);

}

}