OpenCL实现MapReduce算法

版权声明:本文为博主原创文章,未经博主允许不得转载。 https://blog.csdn.net/u011028771/article/details/52624567

MapRedue算法并不是为了解决特定问题而存在的,它是为了解决涉及分布式运算的一类问题而提供的一种计算框架。MapReduce的基本实现包含两部分:映射阶段,从输入数据中产生键值对;归并阶段,处理这些键值对,并输出结果。MapReduce常常和集群计算相关,但是在此仅仅以一个OpenCL实现字符串搜索的例子,对其进行简单介绍。

1.kernel程序解析

kernel.cl
__kernel void mapReduce(char16 pattern,
                        __global char *text,
                        int chars_per_item,
                        __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 item_offset = get_global_id(0)*chars_per_item;

    for (int i = item_offset; i < item_offset + chars_per_item;i++) {
        //text是文本的起始地址;vload函数从text+i地址开始加载16个分量到text_vector中;
        text_vector = vload16(0, text + i);
        check_vector = text_vector == pattern;
        //all(),最高有效位是1,该函数返回1
        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_LOCAL_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]);
    }
}
char16 pattern:字符串"thatwithhavefrom",注意它是私有变量
__global char *text:读入的文本
int chars_per_item:每一个工作项所需要处理的字符数量
__local int * local_result:局部变量,用于存储局部归并结果
__global int* global_result:全局变量,用于存储全局归并结果

pattern作为用于比对的常量,声明为私有变量读取速度比较快;每个工作项都会获得一个pattern副本。
local_result作为局部变量,每个工作组都会维护一个local_result变量;所以对于local_result的初始化需要利用barrier函数同步,确保每一个工作组都初始化完成。
接下来进行数据划分,划分每个工作项处理数据的节点,例如

这里写图片描述
假设现在有96个数据,全局工作项有12个,那么每个工作项需要处理8个数据;那么第一个工作项需要处理0~7的数据,第二个需要处理8~15的数据……
所以

char_per_item=text_size/global_size+1;//(8=95/12+1)

而每个工作项处理的数据区间为:

item_offset~item_offset+chap_per_item;
//(item_offset=0,8,16,……88);

这样global__size个工作项开始同时处理数据;

但是每个工作项处理数据也是采用的并行;vload16()函数一次可以加载16个数据,所以每个工作项,每次可以处理16个数据;
处理结果以工作组为单位归并到local_result中;从上文例图中可以看到,有三个工作组,所以局部归并结果会被存储在三组local_result中。
最后需要进行全局归并,获得最终结果。

2.kernel中工作项间的同步

工作项的同步化只能进行组内同步化,也就是说,只有同一个工作组内的工作项可以同步化,工作组之间不可以同步化。
如果多个工作项访问的是局部内存中的同一段数据,他们之间的执行顺序不同可能得到不同的结果,甚至错误。
为此OpenCL提供了两种同步化方法:

  • 栅栏和障碍
  • 原子操作

2.1 kernel中的原子操作

原子操作的意思是不可分解,例如local_result自加1的运算;在C语言中可以采用如下方式实现:

local_result+=1;

但是kernel程序中不可以这样写,因为local_result是一个工作组公有的空间,假设工作项1读取local_result(值为2),此时刚好工作项2也读取了local_result,这样两个工作项读取到的local_result值是一样的,各自将其加1以后再写入local_result(值为3);但是两个工作项都需要将其加1,最终的结果应该是4才对,但是以为并行操作的原因,出现了错误。如果要得到正确的结果,就需要工作项2等工作项1完成运算以后再进行操作。所以需要采用原子操作来解决这个问题。原子操作是不可被中断的,工作项2不需要等到工作项1运算完成以后再执行,后面的工作项也不需要和工作项2发生关联。
改程序中用到了两个:

atomic_inc()和atomic_add();

2.2kernel中的障碍

 void barrier(cl_mem_fance_flag flags);//障碍
  • CLK_LOCAL_MEM_FENCE:确定内存操作如何影响和工作组局部内存有关的操作
    例如,每个工作组都必须在障碍函数完成同步化之前,处理完局部内存的访问。

  • CL_GLOBAL_MEM_FENCE:同步化全局内存访问

栅栏跟障碍相似,区别在于栅栏主要用来同步特定内存的操作,有些栅栏负责同步化读操作,有些负责同步写操作。

3.主机代码

#include<stdio.h>
#include <stdlib.h>
#include <string.h>

#include<CL/cl.h>
#pragma warning( disable : 4996 )
int main() {
    cl_int error;
    cl_platform_id platforms;

    cl_device_id devices;

    cl_context context;

    FILE *program_handle;
    size_t program_size;
    char *program_buffer;
    cl_program program;

    size_t log_size;
    char *program_log;

    char kernel_name[] = "mapReduce";
    cl_kernel kernel;

    size_t global_size;
    size_t local_size;
    size_t offset = 0;
    cl_command_queue queue;
    //获取平台
    error = clGetPlatformIDs(1, &platforms, NULL);
    if (error != CL_SUCCESS) {
        printf("Get platform failed!");
        return -1;
    }
    error = clGetDeviceIDs(platforms, CL_DEVICE_TYPE_GPU, 1, &devices, NULL);
    if (error != CL_SUCCESS) {
        printf("Get device failed!");
        return -1;
    }
    //获取全局工作项和工作组大小
    clGetDeviceInfo(devices, CL_DEVICE_MAX_COMPUTE_UNITS,
        sizeof(global_size), &global_size, NULL);
    clGetDeviceInfo(devices, CL_DEVICE_MAX_WORK_GROUP_SIZE,
        sizeof(local_size), &local_size, NULL);
    global_size *= local_size;

    //创建上下文
    context = clCreateContext(NULL, 1, &devices, NULL, NULL, &error);
    if (error != CL_SUCCESS) {
        printf("Creat context failed!");
        return -1;
    }
    //创建程序
    program_handle = fopen("kernel.cl", "rb");
    if (program_handle == NULL) {
        printf("The kernle can not be opened!");
        return -1;
    }
    fseek(program_handle, 0, SEEK_END);
    program_size = ftell(program_handle);
    rewind(program_handle);

    program_buffer = (char *)malloc(program_size + 1);
    program_buffer[program_size] = '\0';
    error = fread(program_buffer, sizeof(char), program_size, program_handle);
    if (error == 0) {
        printf("Read kernel failed!");
        return -1;
    }
    fclose(program_handle);
    program = clCreateProgramWithSource(context, 1, (const char **)&program_buffer, &program_size, &error);
    if (error != CL_SUCCESS) {
        printf("Couldn't create the program!");
        return -1;
    }
    //编译程序
    error = clBuildProgram(program, 1, &devices, NULL, NULL, NULL);
    if (error != CL_SUCCESS) {
        //确定日志文件的大小
        clGetProgramBuildInfo(program, devices, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
        program_log = (char *)malloc(log_size + 1);
        program_log[log_size] = '\0';
        //读取日志
        clGetProgramBuildInfo(program, devices, CL_PROGRAM_BUILD_LOG, log_size + 1, program_log, NULL);
        printf("%s\n", program_log);
        free(program_log);
        return -1;
    }
    //创建命令队列
    queue = clCreateCommandQueue(context, devices, CL_QUEUE_PROFILING_ENABLE, &error);
    if (error != CL_SUCCESS) {
        printf("Coudn't create the command queue");
        return -1;
    }
    //创建内核
    kernel = clCreateKernel(program, kernel_name, &error);
    if (kernel == NULL) {
        printf("Couldn't create kernel!\n");
        return -1;
    }
    //初始化参数
    char pattern[16] = "thatwithhavefrom";
    int chars_per_item=0;
    int global_result[4];
    char *textBuffer;
    FILE *tb;
    size_t text_size;
    tb = fopen("textcheck.txt","rt");
    if (tb == NULL) {
        printf("Open textcheck.txt failed!");
        return -1;
    }
    fseek(tb, 0, SEEK_END);
    text_size = ftell(tb)-1;
    fseek(tb, 0, SEEK_SET);
    textBuffer = (char*)malloc(text_size);
    fread(textBuffer, sizeof(char), text_size, tb);
    fclose(tb);
    chars_per_item = text_size /global_size+1;
    //创建缓存对象
    cl_mem memObject1 = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, text_size, textBuffer, &error);
    if (error != CL_SUCCESS) {
        printf("Creat memObject2 failed!\n");
        return -1;
    }
    cl_mem memObject2 = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(int)*4, NULL, &error);
    if (error != CL_SUCCESS) {
        printf("Creat memObject3 failed!\n");
        return -1;
    }
    //设置内核参数
    error = clSetKernelArg(kernel, 0, sizeof(pattern), pattern);
    error |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &memObject1);
    error |= clSetKernelArg(kernel, 2, sizeof(chars_per_item), &chars_per_item);
    error |= clSetKernelArg(kernel, 3, sizeof(int) * 4, NULL);
    error |= clSetKernelArg(kernel, 4, sizeof(cl_mem) , &memObject2);
    if (error != CL_SUCCESS) {
        printf("Error setting kernel arguments!\n");
        return -1;
    }
    //执行内核
    error = clEnqueueNDRangeKernel(queue, kernel, 1, &offset, &global_size, &local_size, 0, NULL, NULL);
    if (error != CL_SUCCESS) {
        printf("Error queuing kernel for execution!\n");
        return -1;
    }
    //读取执行结果
    error = clEnqueueReadBuffer(queue, memObject2, CL_TRUE, 0, sizeof(int)*4,global_result, 0, NULL, NULL);
    if (error != CL_SUCCESS) {
        printf("Error reading result buffer!\n");
        return -1;
    }
    //显示结果
    printf("\nResults: \n");
    printf("Number of occurrences of 'that': %d\n", global_result[0]);
    printf("Number of occurrences of 'with': %d\n", global_result[1]);
    printf("Number of occurrences of 'have': %d\n", global_result[2]);
    printf("Number of occurrences of 'from': %d\n", global_result[3]);

    //释放资源
    clReleaseDevice(devices);
    clReleaseContext(context);
    clReleaseCommandQueue(queue);
    clReleaseProgram(program);
    clReleaseKernel(kernel);
    clReleaseMemObject(memObject1);
    clReleaseMemObject(memObject2);
    return 0;
}

运行结果:
这里写图片描述

没有更多推荐了,返回首页

私密
私密原因:
请选择设置私密原因
  • 广告
  • 抄袭
  • 版权
  • 政治
  • 色情
  • 无意义
  • 其他
其他原因:
120
出错啦
系统繁忙,请稍后再试

关闭