GPGPU OpenCL implements exact string lookup

String lookup is an important operation in the field of information security and information filtering, especially in real-time processing of large text. As an example, the exact pattern string lookup is performed using GPU OpenCL.

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<length; ++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,//p Attern Length const UINT Maxsearchlength//maximum Search positions for each work-group __global int* Ltcount,//result counts (global) __global int* Resultbuffer,//save the match result __local uchar*
    Alpattern)//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<endsearchidx; stringpos+=localsize) {if compare (text+
            Stringpos, Localpattern, patternlength) = = 1) {int count = Atomic_inc (Resultcount);
        Resultbuffer[count] = Stringpos;
        printf ("%d", stringpos);
    } barrier (clk_local_mem_fence); }
}