Mercurial > hg > Game > Cerium

annotate example/many_task/gpu/QuickSort.cl @ 1851:637fa9a4105b draft

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
fix many_task ( sort )
author Shinji KONO <kono@ie.u-ryukyu.ac.jp>
date Sat, 21 Dec 2013 19:42:17 +0900
parents 5c4e3f0d372a
children f800f61a0311
Ignore whitespace changes - Everywhere: Within whitespace: At end of lines:
rev   line source
1520
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
1 typedef struct Data {
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
2 int index;
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
3 int ptr;
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
4 int pad[2];
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
5 } Data, *DataPtr;
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
6
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
7 inline void
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
8 swap(__global Data *data, int left, int right )
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
9 {
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
10 Data tmp = data[left];
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
11 data[left] = data[right];
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
12 data[right] = tmp;
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
13 }
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
14
1538
fac06524090b add gpu task wordcount. But not work print
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents: 1528
diff changeset
15 void
fac06524090b add gpu task wordcount. But not work print
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents: 1528
diff changeset
16 bubble_sort(__global Data *data, int begin, int end)
fac06524090b add gpu task wordcount. But not work print
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents: 1528
diff changeset
17 {
fac06524090b add gpu task wordcount. But not work print
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents: 1528
diff changeset
18 for (int count=0;count<end;count++) {
fac06524090b add gpu task wordcount. But not work print
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents: 1528
diff changeset
19 for (int c=end;c>count;c--) {
fac06524090b add gpu task wordcount. But not work print
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents: 1528
diff changeset
20 if (data[c].index<data[c-1].index)swap(data,c-1,c);
fac06524090b add gpu task wordcount. But not work print
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents: 1528
diff changeset
21 }
fac06524090b add gpu task wordcount. But not work print
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents: 1528
diff changeset
22 }
fac06524090b add gpu task wordcount. But not work print
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents: 1528
diff changeset
23 }
fac06524090b add gpu task wordcount. But not work print
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents: 1528
diff changeset
24
fac06524090b add gpu task wordcount. But not work print
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents: 1528
diff changeset
25
1520
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
26 __kernel void
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
27 quick_sort(__constant int *count,
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
28 __global Data *data)
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
29 {
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
30 int begin = 0;
1528
360b424125ea bug fix gpu kernel file
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents: 1520
diff changeset
31 int end = count[0]-1;
1520
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
32
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
33 int stack[1024];
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
34 int sp = 0;
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
35 int p;
1538
fac06524090b add gpu task wordcount. But not work print
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents: 1528
diff changeset
36
1520
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
37 while (1) {
1544
5c4e3f0d372a many_task add task array
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents: 1542
diff changeset
38 while (begin < end) {
1542
9ccfdc408d51 fix gpu word count.but not count line num.
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents: 1538
diff changeset
39 /*
9ccfdc408d51 fix gpu word count.but not count line num.
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents: 1538
diff changeset
40 * if (end-begin <= 50) {
9ccfdc408d51 fix gpu word count.but not count line num.
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents: 1538
diff changeset
41 * bubble_sort(data,begin,end);
9ccfdc408d51 fix gpu word count.but not count line num.
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents: 1538
diff changeset
42 * break;
9ccfdc408d51 fix gpu word count.but not count line num.
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents: 1538
diff changeset
43 * }
9ccfdc408d51 fix gpu word count.but not count line num.
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents: 1538
diff changeset
44 */
1520
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
45 int where = (begin + end) / 2;
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
46 int pivot = data[where].index;
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
47 data[where].index = data[begin].index;
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
48 int i;
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
49 p = begin;
1851
637fa9a4105b fix many_task ( sort )
Shinji KONO <kono@ie.u-ryukyu.ac.jp>
parents: 1544
diff changeset
50 for (i=begin+1; i<end; i++) {
1520
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
51 if (data[i].index < pivot) {
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
52 p++;
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
53 swap(data, p, i);
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
54 }
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
55 }
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
56 data[begin].index = data[p].index;
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
57 data[p].index = pivot;
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
58
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
59 stack[sp++] = p + 1;
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
60 stack[sp++] = end;
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
61 end = p - 1;
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
62 }
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
63 if (sp == 0) return;
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
64 end = stack[--sp];
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
65 begin = stack[--sp];
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
66 begin = p + 1;
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
67 }
1538
fac06524090b add gpu task wordcount. But not work print
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents: 1528
diff changeset
68
1520
031f26b15ae6 add many_task/gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
69 }