--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/example/word_count/gpu/Exec.cl	Sun Dec 16 21:06:36 2012 +0900
@@ -0,0 +1,36 @@
+__kernel void
+run(__global int *data_count,
+    __global char *i_data,
+    __global int *o_data)
+{
+    __global int *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;
+
+}

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/example/word_count/gpu/Print.cc	Sun Dec 16 21:06:36 2012 +0900
@@ -0,0 +1,92 @@
+#include <stdio.h>
+#include <string.h>
+#include "Print.h"
+#include "Func.h"
+#include "WordCount.h"
+
+/* これは必須 */
+SchedDefineTask1(Print,run_print);
+
+static int
+run_print(SchedTask *s, void *rbuf, void *wbuf)
+{
+    WordCount *w = *(WordCount**)rbuf;
+    unsigned long long *idata = w->o_data;
+    // long task_num = w->task_num;
+    long status_num = w->status_num;
+    int out_task_num = w->out_task_num;
+
+    /*
+     *  head_flag
+     *  o_data[0]
+     *  o_data[1]
+     *
+     */
+    unsigned long long word_data[2];
+
+    int flag_cal_sum = 0;
+    //printf("pad %d\n",pad);
+
+
+    /* head_tail_flag : task = 2 : 1
+     *
+     *    head_tail_flag[2]
+     *  ___________ __________
+     * |           |          |
+     * | head_flag | tail_fag | 
+     * |___________|__________|
+     *      [0]         [1] 
+     * 
+     * |----------------------|
+     *    unsigned long long 
+     *          16byte
+     *
+     * 配列先頭の head_flag はみない
+     * 配列末尾の tail_flag はみない
+     * 担当範囲前の末尾文字が「改行と、スペース以外」(tail_flag = 1)で、
+     * 担当範囲の先頭文字が「改行とスペース」(head_flag = 0)の場合
+     * 単語数を +1  することで、整合性がとれる。
+     *
+     * ex.
+     *    task_num 4
+     *    head_tail_flag[8]
+     *
+     *      0   1   2   3   4   5   6   7
+     *     ___ ___ ___ ___ ___ ___ ___ ___ 
+     *    | h | t | h | t | h | t | h | t |
+     *    |___|___|___|___|___|___|___|___|
+     *        |-------|-------|-------| 
+     *           比較    比較    比較
+     *
+     */
+
+    s->printf("start sum\n");
+
+    for (int i = 0; i < status_num; i++) {
+	word_data[i] = 0;
+    }
+
+    for (int i = 0; i < out_task_num ; i++) {
+	word_data[0] += idata[i*w->out_size+0];
+	word_data[1] += idata[i*w->out_size+1];
+	unsigned long long *head_tail_flag = 
+		       &idata[i*w->out_size+2];
+        if((i!=out_task_num-1)&&
+           (head_tail_flag[1] == 1) && (head_tail_flag[4] == 0)) {
+	    flag_cal_sum++;
+        }
+    }
+
+    word_data[0] += flag_cal_sum;
+
+    for (int i = status_num-1; i >=0; i--) {
+	s->printf("%llu ",word_data[i]);
+    }
+
+    //printf("j = %d\n",j);
+
+
+    s->printf("\n");
+
+    return 0;
+}

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/example/word_count/gpu/Print.h	Sun Dec 16 21:06:36 2012 +0900
@@ -0,0 +1,9 @@
+#ifndef INCLUDED_TASK_PRINT
+#define INCLUDED_TASK_PRINT
+
+#ifndef INCLUDED_SCHED_TASK
+#  include "SchedTask.h"
+#endif
+
+
+#endif

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/example/word_count/gpu/task_init.cc	Sun Dec 16 21:06:36 2012 +0900
@@ -0,0 +1,20 @@
+#include "Func.h"
+#include "Scheduler.h"
+#include "GpuScheduler.h"
+
+/* 必ずこの位置に書いて */
+SchedExternTask(Print);
+SchedExternTask(RUN_TASK_BLOCKS);
+
+/**
+ * この関数は ../spe/spe-main と違って
+ * 自分で呼び出せばいい関数なので
+ * 好きな関数名でおk (SchedRegisterTask は必須)
+ */
+void
+task_init(void)
+{
+    SchedRegisterTask(TASK_PRINT, Print);
+    GpuSchedRegister(TASK_EXEC, "gpu/Exec.cl", "run");
+    SchedRegister(RUN_TASK_BLOCKS);
+}