Mercurial > hg > Game > Cerium

changeset 1837:093bdd5e940e draft

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
add file
author Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
date Fri, 20 Dec 2013 12:00:15 +0900
parents 56692133c5fb
children 759587e37bc7
files example/word_count/gpu/Exec.cl example/word_count/gpu/Exec_Data_Parallel.cl example/word_count/main.cc
diffstat 3 files changed, 46 insertions(+), 2 deletions(-) [+]
line wrap: on
line diff
--- a/example/word_count/gpu/Exec.cl	Fri Dec 20 04:36:27 2013 +0900
+++ b/example/word_count/gpu/Exec.cl	Fri Dec 20 12:00:15 2013 +0900
@@ -3,7 +3,7 @@
           __global char *i_data,
           __global unsigned long long *o_data)
 {
-    __global unsigned long long *head_tail_flag = o_data +2;
+    __global unsigned long long *head_tail_flag = o_data+2;
     long length = param[0];
     int word_flag = 0;
     int word_num = 0;
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/example/word_count/gpu/Exec_Data_Parallel.cl	Fri Dec 20 12:00:15 2013 +0900
@@ -0,0 +1,44 @@
+__kernel void
+wordcount_parallel(__constant long *param,
+                   __global char *rbuf,
+                   __global unsigned long long *wbuf)
+{
+    long task_spwaned = (long)param[0];
+    long division_size = (long)param[1];
+    long length = (long)param[2];
+    long out_size = (long)param[3];
+    
+    long allocation = task_spwaned + (long)get_global_id(0);
+    
+    __global char *i_data = rbuf + allocation*division_size;
+    __global unsigned long long *o_data = wbuf + allocation*out_size;
+    __global unsigned long long *head_tail_flag = o_data +2;
+    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;
+
+}
--- a/example/word_count/main.cc	Fri Dec 20 04:36:27 2013 +0900
+++ b/example/word_count/main.cc	Fri Dec 20 12:00:15 2013 +0900
@@ -284,7 +284,7 @@
                                    (memaddr)&w->self,sizeof(memaddr),0,0);
     w->t_print = t_print;
     
-    for(int i=0;i<4;i++) {
+    for(int i=0;i<1;i++) {
         /* Task を task_blocks ずつ起動する Task */
         /* serialize されていると仮定する... */
         HTaskPtr t_exec = manager->create_task(RUN_TASK_BLOCKS,