annotate WordCount/word_count.cl @ 2:1a32564347d5

parallel word count
author Yutaka_Kinjyo
date Sun, 24 Jul 2011 01:45:40 +0900
parents 403e35dd9b6d
children
Ignore whitespace changes - Everywhere: Within whitespace: At end of lines:
rev   line source
1
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
1 typedef struct {
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
2 int work_num;
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
3 int size;
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
4 int remain_size;
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
5 } param_t;
0
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
6
1
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
7 __kernel void word_count(__global char* text_data, __global param_t* param, __global int* out_data)
0
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
8 {
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
9
1
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
10 /*
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
11 * out_data
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
12 *
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
13 * ___int____ ___int____ ___int____ ___int____
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
14 * __________ __________ __________ __________
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
15 * | | | | |
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
16 * | word_num | line_num | head | tail |
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
17 * |__________|__________|__________|__________|
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
18 *
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
19 */
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
20
0
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
21
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
22 uint gid = get_global_id(0);
1
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
23 int index = gid*4;
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
24 int head = gid*4+2;
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
25 int tail = gid*4+3;
0
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
26
1
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
27 int work_num = param->work_num;
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
28 int size = param->size;
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
29 int remain_size = param->remain_size;
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
30 // gid が最後なら remain_size を見る。gid は 0 からスタート
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
31 int length = ( (work_num - 1) != gid ) ? size : remain_size;
0
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
32 int word_flag = 0;
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
33 int word_num = 0;
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
34 int line_num = 0;
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
35 int i = 0;
1
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
36 int start = gid*size;
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
37 int end = start + length;
0
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
38
1
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
39 out_data[head] = (text_data[0] != 0x20) && (text_data[0] != 0x0A);
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
40 word_num -= 1-out_data[head];
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
41
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
42
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
43 //for (i = 0; i < length; i++) {
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
44 for (i = start; i < end; i++) {
0
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
45 if (text_data[i] == 0x20) {
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
46 word_flag = 1;
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
47 } else if (text_data[i] == 0x0A) {
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
48 line_num += 1;
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
49 word_flag = 1;
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
50 } else {
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
51 word_num += word_flag;
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
52 word_flag = 0;
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
53 }
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
54 }
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
55
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
56 word_num += word_flag;
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
57
1
403e35dd9b6d word count ...
Yutaka_Kinjyo
parents: 0
diff changeset
58 out_data[tail] = (text_data[i-1] != 0x20) && (text_data[i-1] != 0x0A);
0
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
59
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
60 out_data[index] = word_num;
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
61 out_data[index+1] = line_num;
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
62
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
63 return 0;
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
64
0e6e76dbdb0f add file
Yutaka_Kinjyo
parents:
diff changeset
65 }