Mercurial > hg > Game > Cerium

comparison example/word_count/cuda/Exec_Data_Parallel.cu @ 1941:f19885ea776d draft

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
add wordcount for cuda. fix CudaScheduler. add makefile
author Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
date Tue, 04 Feb 2014 02:18:07 +0900
parents
children
comparison
equal deleted inserted replaced
1936:e8ca9cae59fc 1941:f19885ea776d
1 extern "C" {
2 __global__ void
3 wordcount_parallel(long *param,
4 char *rbuf,
5 unsigned long *wbuf)
6 {
7 long task_spwaned = param[0];
8 long division_size = param[1];
9 long length = param[2];
10 long out_size = param[3];
11 int allocation = (int)task_spwaned + (blockIdx.x * blockDim.x + threadIdx.x);
12 char *i_data = rbuf + allocation*division_size;
13 unsigned long *o_data = wbuf + allocation*out_size;
14 unsigned long *head_tail_flag = o_data+2;
15 long word_flag = 0;
16 long word_num = 0;
17 long line_num = 0;
18 long i = 0;
19
20 head_tail_flag[0] = (i_data[0] != 0x20) && (i_data[0] != 0x0A);
21 word_num -= 1-head_tail_flag[0];
22
23 for (; i < length; i++) {
24 if (i_data[i] == 0x20) {
25 word_flag = 1;
26 } else if (i_data[i] == 0x0A) {
27 line_num += 1;
28 word_flag = 1;
29 } else {
30 word_num += word_flag;
31 word_flag = 0;
32 }
33 }
34
35 word_num += word_flag;
36 head_tail_flag[1] = (i_data[i-1] != 0x20) && (i_data[i-1] != 0x0A);
37
38 o_data[0] = (unsigned long)word_num;
39 o_data[1] = (unsigned long)line_num;
40 }
41 }