GPGPU OpenCL實現精確字串尋找

字串尋找是資訊安全、資訊過濾領域的重要操作，尤其是對大文本的即時處理。這篇作為執行個體，使用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)：

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