__kernel void
wordcount(__constant long *param,
          __global char *i_data,
          __global unsigned long *o_data)
{
    __global unsigned long *head_tail_flag = o_data+2;
    long length = param[0];
    long word_flag = 0;
    long word_num = 0;
    long line_num = 0;
    long 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)word_num;
    o_data[1] = (unsigned long)line_num;

}