__kernel void
wordcount(__global int *data_count,
    __global void *rbuf,
    __global void *wbuf)
{
    __global char *i_data =  (__global char *)rbuf;
    __global unsigned long long *o_data = (__global unsigned long long*)wbuf;
    __global  unsigned long long *head_tail_flag = o_data +2;
    int length = data_count[0];
    int word_flag = 0;
    int word_num = 0;
    int line_num = 0;
    int i = 0;

    head_tail_flag[0] = (i_data[0] != 0x20) && (i_data[0] != 0x0A);
    word_num -= 1-head_tail_flag[0];

    for (; i < length; i++) {
        if (i_data[i] == 0x20) {
            word_flag = 1;
        } else if (i_data[i] == 0x0A) {
            line_num += 1;
            word_flag = 1;
        } else {
            word_num += word_flag;
            word_flag = 0;
        }
    }

    word_num += word_flag;
    head_tail_flag[1] = (i_data[i-1] != 0x20) && (i_data[i-1] != 0x0A);

    // s->printf("SPE word %d line %d\n",word_num,line_num);

    o_data[0] = (unsigned long long)word_num;
    o_data[1] = (unsigned long long)line_num;

}