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