GPGPU OpenCL implements exact string lookup

Source: Internet
Author: User
Tags count printf thread

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 (

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); }

Contact Us

The content source of this page is from Internet, which doesn't represent Alibaba Cloud's opinion; products and services mentioned on that page don't have any relationship with Alibaba Cloud. If the content of the page makes you feel confusing, please write us an email, we will handle the problem within 5 days after receiving your email.

If you find any instances of plagiarism from the community, please send an email to: and provide relevant evidence. A staff member will contact you within 5 working days.

A Free Trial That Lets You Build Big!

Start building with 50+ products and up to 12 months usage for Elastic Compute Service

  • Sales Support

    1 on 1 presale consultation

  • After-Sales Support

    24/7 Technical Support 6 Free Tickets per Quarter Faster Response

  • Alibaba Cloud offers highly flexible support services tailored to meet your exact needs.