OpenCL实现MapReduce算法

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函数同步，确保每一个工作组都初始化完成。
接下来进行数据划