精確字串搜尋在資訊安全、日誌分析、基因序列比對等領域極為常見。面對 GB 級文本，CPU 串行掃描效率低下。GPU 的大量執行緒可將文本分塊並行處理，但需解決模式串共享、邊界條件、結果寫入衝突三大挑戰。
| 記憶體類型 | 存放內容 | 優勢 |
|---|---|---|
| Private Memory | 模式串長度、文本長度、搜尋範圍 | 零延遲、每執行緒獨有 |
| Local Memory | 模式串內容 | 工作組內共享、低延遲、減少全域頻寬 |
| Global Memory | 待搜尋文本、結果陣列 | 容量大、全裝置可見 |
| Atomic Counter | 結果計數器 | 無鎖並行計數、保證寫入位置唯一 |
// 輔助函數：比較模式串
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, // 輸入文本
const uint text_length, // 文本長度
__global const uchar* pattern, // 模式串
const uint pattern_length, // 模式串長度
const uint max_search_length, // 每工作組最大搜尋位置數
__global int* result_count, // 結果計數器 (原子操作)
__global int* result_buffer, // 結果緩衝區
__local uchar* local_pattern // 本地記憶體模式串緩衝
) {
int local_id = get_local_id(0);
int local_size = get_local_size(0);
int group_id = get_group_id(0);
// 計算全域有效搜尋範圍
uint last_valid_pos = text_length - pattern_length + 1;
uint start_pos = group_id * max_search_length;
uint end_pos = start_pos + max_search_length;
// 邊界檢查
if (start_pos > last_valid_pos) return;
if (end_pos > last_valid_pos) end_pos = last_valid_pos;
// 階段 1：協作載入模式串到本地記憶體
for (int idx = local_id; idx < pattern_length; idx += local_size) {
local_pattern[idx] = pattern[idx];
}
barrier(CLK_LOCAL_MEM_FENCE); // 確保所有 work-item 看到完整模式串
// 階段 2：並行搜尋
for (uint pos = start_pos + local_id; pos < end_pos; pos += local_size) {
if (compare(text + pos, local_pattern, pattern_length)) {
// 原子遞增獲取唯一寫入索引
int count = atomic_inc(result_count);
result_buffer[count] = pos;
}
// 注意：此處 barrier 僅為示範，實際搜尋迴路不需每輪同步
// 除非有共享資料相依，否則移除可提升效能
}
}
#include <CL/cl.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#define CHECK_ERR(err, msg) \
if (err != CL_SUCCESS) { fprintf(stderr, "%s: %d\n", msg, err); exit(1); }
int main() {
cl_int err;
// 1. 平台與設備初始化
cl_platform_id platform;
clGetPlatformIDs(1, &platform, NULL);
cl_device_id device;
clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
cl_context ctx = clCreateContext(NULL, 1, &device, NULL, NULL, &err);
CHECK_ERR(err, "Context");
cl_command_queue queue = clCreateCommandQueue(ctx, device, 0, &err);
CHECK_ERR(err, "Queue");
// 2. 載入 Kernel 原始碼
FILE* fp = fopen("string_search_kernel.cl", "r");
fseek(fp, 0, SEEK_END);
size_t src_len = ftell(fp);
rewind(fp);
char* src = (char*)malloc(src_len + 1);
fread(src, 1, src_len, fp);
src[src_len] = '\0';
fclose(fp);
cl_program prog = clCreateProgramWithSource(ctx, 1, (const char**)&src, &src_len, &err);
CHECK_ERR(err, "Program");
err = clBuildProgram(prog, 1, &device, NULL, NULL, NULL);
if (err != CL_SUCCESS) {
size_t log_size;
clGetProgramBuildInfo(prog, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
char* log = (char*)malloc(log_size);
clGetProgramBuildInfo(prog, device, CL_PROGRAM_BUILD_LOG, log_size, log, NULL);
fprintf(stderr, "Build failed:\n%s\n", log);
free(log);
exit(1);
}
free(src);
cl_kernel kernel = clCreateKernel(prog, "StringSearch", &err);
CHECK_ERR(err, "Kernel");
// 3. 準備輸入資料
// 讀取文本檔案
FILE* tfp = fopen("StringSearch_Input.txt", "rb");
fseek(tfp, 0, SEEK_END);
size_t text_len = ftell(tfp);
rewind(tfp);
char* text = (char*)malloc(text_len + 1);
fread(text, 1, text_len, tfp);
text[text_len] = '\0';
fclose(tfp);
const char* pattern = "info";
int pattern_len = strlen(pattern);
int max_search_len = 256 * 64; // 每工作組處理 16K 字元
// 結果緩衝區
int* result_count = (int*)calloc(1, sizeof(int));
int* results = (int*)calloc(text_len, sizeof(int));
// 4. 建立 Buffer
cl_mem text_buf = clCreateBuffer(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
text_len, text, &err);
CHECK_ERR(err, "Text buffer");
cl_mem pattern_buf = clCreateBuffer(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
pattern_len, pattern, &err);
CHECK_ERR(err, "Pattern buffer");
cl_mem count_buf = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
sizeof(int), result_count, &err);
CHECK_ERR(err, "Count buffer");
cl_mem result_buf = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
text_len * sizeof(int), results, &err);
CHECK_ERR(err, "Result buffer");
// 5. 設定 Kernel 參數
clSetKernelArg(kernel, 0, sizeof(cl_mem), &text_buf);
clSetKernelArg(kernel, 1, sizeof(int), &text_len);
clSetKernelArg(kernel, 2, sizeof(cl_mem), &pattern_buf);
clSetKernelArg(kernel, 3, sizeof(int), &pattern_len);
clSetKernelArg(kernel, 4, sizeof(int), &max_search_len);
clSetKernelArg(kernel, 5, sizeof(cl_mem), &count_buf);
clSetKernelArg(kernel, 6, sizeof(cl_mem), &result_buf);
// 第 7 參數：本地記憶體大小 = pattern_len
clSetKernelArg(kernel, 7, pattern_len, NULL);
// 6. 計算工作維度
int global_work_items = (text_len + 63) / 64; // 每工作組 256 threads，每 thread 處理 64 字元
size_t global_ws = global_work_items;
size_t local_ws = 256;
// 7. 執行
cl_event event;
err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_ws, &local_ws, 0, NULL, &event);
CHECK_ERR(err, "EnqueueKernel");
clWaitForEvents(1, &event);
clReleaseEvent(event);
// 8. 讀回結果
clEnqueueReadBuffer(queue, count_buf, CL_TRUE, 0, sizeof(int), result_count, 0, NULL, NULL);
printf("Matches found: %d\n", *result_count);
if (*result_count > 0) {
clEnqueueReadBuffer(queue, result_buf, CL_TRUE, 0,
*result_count * sizeof(int), results, 0, NULL, NULL);
for (int i = 0; i < *result_count; i++) {
printf("Position %d: %d\n", i, results[i]);
}
}
// 9. 清理
clReleaseMemObject(text_buf);
clReleaseMemObject(pattern_buf);
clReleaseMemObject(count_buf);
clReleaseMemObject(result_buf);
clReleaseKernel(kernel);
clReleaseProgram(prog);
clReleaseCommandQueue(queue);
clReleaseContext(ctx);
free(text);
free(pattern);
free(result_count);
free(results);
return 0;
}
pos += local_size 確保 work-item 訪問連續記憶體區域，實現記憶體合併訪問atomic_inc 保證多執行緒寫入結果時索引唯一，無需鎖
