--- a/WordCount/word_count.cl	Tue Jul 12 11:12:51 2011 +0900
+++ b/WordCount/word_count.cl	Tue Jul 12 17:17:11 2011 +0900
@@ -1,26 +1,47 @@
-#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable
+typedef struct {
+    int work_num;
+    int size;
+    int remain_size;
+} param_t;
 
-__kernel void word_count(__global char* text_data,  __global int* size, __global int* out_data)
+__kernel void word_count(__global char* text_data,  __global param_t* param, __global int* out_data)
 {
 
-  //unsigned long long *head_tail_flag = o_data +2;
+    /*
+     *                   out_data  
+     *
+     *  ___int____ ___int____ ___int____ ___int____ 
+     *  __________ __________ __________ __________ 
+     * |          |          |          |          |
+     * | word_num | line_num |   head   |   tail   |
+     * |__________|__________|__________|__________|
+     *
+     */
+
 
   uint gid = get_global_id(0);
+  int index = gid*4;
+  int head = gid*4+2;
+  int tail = gid*4+3;
 
-  int length = *size;
+  int work_num = param->work_num;
+  int size = param->size;
+  int remain_size = param->remain_size;
+  // gid が最後なら remain_size を見る。gid は 0 からスタート
+  int length = ( (work_num - 1) != gid ) ? size : remain_size;
   int word_flag = 0;
   int word_num = 0;
   int line_num = 0;
   int i = 0;
-    
-  int start = gid*length/2;
-  int end = start + length/2;
-  
-  //head_tail_flag[0] = (i_data[0] != 0x20) && (i_data[0] != 0x0A);
-  //word_num -= 1-head_tail_flag[0];
+  int start = gid*size;
+  int end = start + length;
 
-  for (i = 0; i < length; i++) {
-  //for (i = start; i < end; i++) {
+  out_data[head] = (text_data[0] != 0x20) && (text_data[0] != 0x0A);
+  word_num -= 1-out_data[head];
+
+
+  //for (i = 0; i < length; i++) {
+  for (i = start; i < end; i++) {
     if (text_data[i] == 0x20) {
       word_flag = 1;
     } else if (text_data[i] == 0x0A) {
@@ -34,17 +55,11 @@
 
   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);
+  out_data[tail] = (text_data[i-1] != 0x20) && (text_data[i-1] != 0x0A);
 
-  int index = gid*2;
-  
   out_data[index] = word_num;
   out_data[index+1] = line_num;
 
-
-  
   return 0;
 
 }