extern "C" {
    __global__ void
    wordcount(long *param,
              char *i_data,
              unsigned long *o_data)
    {
        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);
        
        o_data[0] = (unsigned long)word_num;
        o_data[1] = (unsigned long)line_num;
    }
}