GPGPU OpenCL How to achieve accurate string lookup?

Source: Internet
Author: User
Tags printf

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)

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 ENDSEARCHIDX = Beginsearchidx + maxsearchlength;

if (Beginsearchidx > Lastsearchidx)


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



Related Article

E-Commerce Solutions

Leverage the same tools powering the Alibaba Ecosystem

Learn more >

Apsara Conference 2019

The Rise of Data Intelligence, September 25th - 27th, Hangzhou, China

Learn more >

Alibaba Cloud Free Trial

Learn and experience the power of Alibaba Cloud with a free trial worth $300-1200 USD

Learn more >

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.