|
1673
|
1 __kernel
|
|
|
2 void butterfly(__constant long *param, __global float2 *x_in, __global float2 *w, __global float2 *x_out)
|
|
|
3 {
|
|
|
4 unsigned long gid = (unsigned long)get_global_id(0);
|
|
|
5 unsigned long nid = (unsigned long)get_global_id(1);
|
|
|
6
|
|
1835
|
7 long n = param[0];
|
|
|
8 unsigned long direction_flag = (unsigned long)param[1];
|
|
|
9 long iter = param[2];
|
|
1673
|
10
|
|
|
11 int butterflySize = 1 << (iter-1);
|
|
|
12 int butterflyGrpDist = 1 << iter;
|
|
|
13 int butterflyGrpNum = n >> iter;
|
|
|
14 int butterflyGrpBase = (gid >> (iter-1))*(butterflyGrpDist);
|
|
|
15 int butterflyGrpOffset = gid & (butterflySize-1);
|
|
|
16
|
|
|
17 int a = nid * n + butterflyGrpBase + butterflyGrpOffset;
|
|
|
18 int b = a + butterflySize;
|
|
|
19
|
|
|
20 int l = butterflyGrpNum * butterflyGrpOffset;
|
|
|
21
|
|
|
22 float2 xa, xb, xbxx, xbyy, wab, wayx, wbyx, resa, resb;
|
|
|
23
|
|
|
24 xa = x_in[a];
|
|
|
25 xb = x_in[b];
|
|
|
26 xbxx.x = xbxx.y = xb.x;
|
|
|
27 xbyy.x = xbyy.y = xb.y;
|
|
|
28
|
|
|
29 wab.x = w[l].x;
|
|
|
30 if(direction_flag == 0x80000000) {
|
|
|
31 wab.y = -w[l].y;
|
|
|
32 } else {
|
|
|
33 wab.y = w[l].y;
|
|
|
34 }
|
|
|
35
|
|
|
36 wayx.x = -wab.y;
|
|
|
37 wayx.y = wab.x;
|
|
|
38
|
|
|
39 wbyx.x = wab.y;
|
|
|
40 wbyx.y = -wab.x;
|
|
|
41
|
|
|
42 resa.x = xa.x + xbxx.x*wab.x + xbyy.x*wayx.x;
|
|
|
43 resa.y = xa.y + xbxx.y*wab.y + xbyy.y*wayx.y;
|
|
|
44
|
|
|
45 resb.x = xa.x - xbxx.x*wab.x + xbyy.x*wbyx.x;
|
|
|
46 resb.y = xa.y - xbxx.y*wab.y + xbyy.y*wbyx.y;
|
|
|
47
|
|
|
48 x_out[a] = resa;
|
|
|
49 x_out[b] = resb;
|
|
|
50 }
|
