Mercurial > hg > Members > yutaka > OpenCL

comparison WordCount/word_count.cl @ 1:403e35dd9b6d

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
word count ...
author Yutaka_Kinjyo
date Tue, 12 Jul 2011 17:17:11 +0900
parents 0e6e76dbdb0f
children
comparison
equal deleted inserted replaced
0:0e6e76dbdb0f 1:403e35dd9b6d
1 #pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable 1 typedef struct {
2 int work_num;
3 int size;
4 int remain_size;
5 } param_t;
2 6
3 __kernel void word_count(__global char* text_data, __global int* size, __global int* out_data) 7 __kernel void word_count(__global char* text_data, __global param_t* param, __global int* out_data)
4 { 8 {
5 9
6 //unsigned long long *head_tail_flag = o_data +2; 10 /*
11 * out_data
12 *
13 * ___int____ ___int____ ___int____ ___int____
14 * __________ __________ __________ __________
15 * | | | | |
16 * | word_num | line_num | head | tail |
17 * |__________|__________|__________|__________|
18 *
19 */
20
7 21
8 uint gid = get_global_id(0); 22 uint gid = get_global_id(0);
23 int index = gid*4;
24 int head = gid*4+2;
25 int tail = gid*4+3;
9 26
10 int length = *size; 27 int work_num = param->work_num;
28 int size = param->size;
29 int remain_size = param->remain_size;
30 // gid が最後なら remain_size を見る。gid は 0 からスタート
31 int length = ( (work_num - 1) != gid ) ? size : remain_size;
11 int word_flag = 0; 32 int word_flag = 0;
12 int word_num = 0; 33 int word_num = 0;
13 int line_num = 0; 34 int line_num = 0;
14 int i = 0; 35 int i = 0;
15 36 int start = gid*size;
16 int start = gid*length/2; 37 int end = start + length;
17 int end = start + length/2;
18
19 //head_tail_flag[0] = (i_data[0] != 0x20) && (i_data[0] != 0x0A);
20 //word_num -= 1-head_tail_flag[0];
21 38
22 for (i = 0; i < length; i++) { 39 out_data[head] = (text_data[0] != 0x20) && (text_data[0] != 0x0A);
23 //for (i = start; i < end; i++) { 40 word_num -= 1-out_data[head];
41
42
43 //for (i = 0; i < length; i++) {
44 for (i = start; i < end; i++) {
24 if (text_data[i] == 0x20) { 45 if (text_data[i] == 0x20) {
25 word_flag = 1; 46 word_flag = 1;
26 } else if (text_data[i] == 0x0A) { 47 } else if (text_data[i] == 0x0A) {
27 line_num += 1; 48 line_num += 1;
28 word_flag = 1; 49 word_flag = 1;
32 } 53 }
33 } 54 }
34 55
35 word_num += word_flag; 56 word_num += word_flag;
36 57
37 //head_tail_flag[1] = (i_data[i-1] != 0x20) && (i_data[i-1] != 0x0A); 58 out_data[tail] = (text_data[i-1] != 0x20) && (text_data[i-1] != 0x0A);
38
39 // s->printf("SPE word %d line %d\n",word_num,line_num);
40 59
41 int index = gid*2;
42
43 out_data[index] = word_num; 60 out_data[index] = word_num;
44 out_data[index+1] = line_num; 61 out_data[index+1] = line_num;
45 62
46
47
48 return 0; 63 return 0;
49 64
50 } 65 }