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

Util.cpp

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

main.cpp

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

test.cl

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