OpenCL-- 统计给定单词在文本中出现次数

#include "util.h"
#include <stdio.h>

cl_int init (cl_platform_id &platform, cl_device_id &device, cl_context &context, cl_command_queue &queue)
{
cl_int err = clGetPlatformIDs (1, &platform, NULL);
if (err < 0){
perror ("Couldn't identify a platform");
return err;
}
err = clGetDeviceIDs (platform, CL_DEVICE_TYPE_GPU, 1, &device,NULL);
if (err < 0){
perror ("Couldn't access any devices");
return err;
}
context = clCreateContext (NULL, 1, &device, NULL, NULL, &err);
if (err < 0){
perror ("Couldn't create a context");
return err;
}
queue = clCreateCommandQueue (context, device, 0, &err);
if (err < 0){
clReleaseContext (context);
perror ("Couldn't create a command queue");
return err;
}
return 0;
}

cl_int create_program (cl_context &context, cl_device_id &device, const char* program_file_name, cl_program &program){
FILE *file = fopen (program_file_name, "r");
cl_int err;
fseek (file, 0, SEEK_END);
size_t program_size = ftell (file);
rewind (file);
char *program_buffer = (char *) malloc (program_size + 1);
fread (program_buffer, sizeof(char), program_size, file);
fclose (file);

program = clCreateProgramWithSource (context, 1, (const char **)&program_buffer, &program_size, &err);
if (err < 0){
perror ("Couldn't create the program");
return err;
}
err = clBuildProgram (program, 1, &device, NULL, NULL, NULL);
if (err < 0){
clReleaseProgram (program);
perror ("Couldn't build the program");
return err;
}
return 0;
}

cl_int create_kernel (cl_program &program, const char *kernel_fun, cl_kernel &kernel){
cl_int err;
kernel = clCreateKernel (program, kernel_fun, &err);
if (err < 0){
perror ("Couldn't create a kernel");
return err;
}
return 0;
}

#include "util.h"
#include <stdio.h>

#define PROGRAM_FILE "test.cl"
#define KERNEL_NAME "vector_add"
#define TEXT_FILE "kafka.txt"

cl_int run_kernel (cl_command_queue &queue, cl_context &context, cl_kernel &kernel)
{
cl_int err;

FILE *file = fopen (TEXT_FILE, "r");
fseek (file, 0, SEEK_END);
size_t text_size = ftell (file);
rewind (file);
char *text_buffer = (char *)malloc (text_size + 1);
fread (text_buffer, sizeof(char), text_size, file);
fclose (file);

int result[4] = {0,0,0,0};

cl_mem a_mem_obj = clCreateBuffer (context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, text_size * sizeof(char), text_buffer, &err);
if (err < 0){
perror ("create buffer error");
return err;
}
free (text_buffer);

cl_mem b_mem_obj = clCreateBuffer (context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, 4 * sizeof(int), result, &err);
if (err < 0){
clReleaseMemObject (a_mem_obj);
perror ("create buffer error");
return err;
}

size_t global_size[1] = {100};
size_t local_size[1] = {10};
char pattern[17] = "thatwithhavefrom";

int chars_per_line = (int)text_size / (int)global_size[0] + 1;

err = clSetKernelArg (kernel, 0, sizeof (cl_mem), &a_mem_obj);
err = clSetKernelArg (kernel, 1, 16, pattern);
err = clSetKernelArg (kernel, 2, sizeof (int), &chars_per_line);
err = clSetKernelArg (kernel, 3, sizeof (int) * 4, NULL);
err = clSetKernelArg (kernel, 4, sizeof(cl_mem), &b_mem_obj);

err = clEnqueueNDRangeKernel (queue, kernel, 1, NULL, global_size, local_size, 0, NULL, NULL);

cl_int *b = (cl_int *)clEnqueueMapBuffer (queue, b_mem_obj, CL_TRUE, CL_MAP_READ, 0, 4 * sizeof(int), 0, NULL, NULL, NULL);
for (int i = 0; i < 4; i++){
printf ("%d\n", b[i]);
}
clEnqueueUnmapMemObject (queue, b_mem_obj, b, 0, NULL, NULL);

clReleaseMemObject (a_mem_obj);
clReleaseMemObject (b_mem_obj);
return 0;
}

int main (){
cl_platform_id platform;
cl_device_id device;
cl_context context;
cl_command_queue queue;
cl_program program;
cl_kernel kernel;
cl_int err;

err = init (platform, device, context, queue);
if (err < 0){
return 1;
}

err = create_program (context, device, PROGRAM_FILE, program);
if (err < 0){
clReleaseCommandQueue (queue);
clReleaseContext (context);
return 1;
}

err = create_kernel (program, KERNEL_NAME, kernel);
if (err < 0){
clReleaseCommandQueue (queue);
clReleaseProgram (program);
clReleaseContext (context);
return 1;
}

run_kernel (queue, context, kernel);

clReleaseKernel (kernel);
clReleaseCommandQueue (queue);
clReleaseProgram (program);
clReleaseContext (context);
return 0;
}

__kernel void vector_add (__global char *text, char16 pattern, int chars_per_line,
__local int *local_result, __global int *global_result){

char16 text_vector, check_vector;

local_result[0] = 0;
local_result[1] = 0;
local_result[2] = 0;
local_result[3] = 0;

barrier (CLK_LOCAL_MEM_FENCE);

int offset = chars_per_line * get_global_id (0);

for (int i = offset; i < offset + chars_per_line; i++){

text_vector = vload16 (0, text + i);

check_vector = text_vector == pattern;

if (all (check_vector.s0123))
atomic_inc (local_result);
if (all (check_vector.s4567))
atomic_inc (local_result + 1);
if (all (check_vector.s89AB))
atomic_inc (local_result + 2);
if (all (check_vector.sCDEF))
atomic_inc (local_result + 3);
}

barrier (CLK_GLOBAL_MEM_FENCE);

if (get_local_id (0) == 0){
atomic_add (global_result, local_result[0]);
atomic_add (global_result + 1, local_result[1]);
atomic_add (global_result + 2, local_result[2]);
atomic_add (global_result + 3, local_result[3]);
}
}