-
Notifications
You must be signed in to change notification settings - Fork 0
/
Copy pathrowTopo16.c
112 lines (106 loc) · 2.38 KB
/
rowTopo16.c
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
#include "coreId16.inc"
void initRing(unsigned int * RIGHT,unsigned int * LEFT, unsigned int * ringOrder, unsigned int * gidOrder)
{
unsigned int gid = get_global_id(0);
unsigned int coreId = LOCAL_MEM_ADDRESS_BASE(gid);
gidOrder[0] = 0;
gidOrder[1] = 1;
gidOrder[2] = 2;
gidOrder[3] = 3;
gidOrder[4] = 7;
gidOrder[5] = 6;
gidOrder[6] = 5;
gidOrder[7] = 4;
gidOrder[8] = 8;
gidOrder[9] = 9;
gidOrder[10] = 10;
gidOrder[11] = 11;
gidOrder[12] = 15;
gidOrder[13] = 14;
gidOrder[14] = 13;
gidOrder[15] = 12;
switch(coreId)
{
case core00:
(*RIGHT) = core10;
(*LEFT) = 0x0;
(*ringOrder) = 0;
break;
case core10:
(*RIGHT) = core20;
(*LEFT) = core00;
(*ringOrder) = 1;
break;
case core20:
(*RIGHT) = core30;
(*LEFT) = core10;
(*ringOrder) = 2;
break;
case core30:
(*RIGHT) = core31;
(*LEFT) = core20;
(*ringOrder) = 3;
break;
case core31:
(*RIGHT) = core21;
(*LEFT) = core30;
(*ringOrder) = 4;
break;
case core21:
(*RIGHT) = core11;
(*LEFT) = core31;
(*ringOrder) = 5;
break;
case core11:
(*RIGHT) = core01;
(*LEFT) = core21;
(*ringOrder) = 6;
break;
case core01:
(*RIGHT) = core02;
(*LEFT) = core11;
(*ringOrder) = 7;
break;
case core02:
(*RIGHT) = core12;
(*LEFT) = core01;
(*ringOrder) = 8;
break;
case core12:
(*RIGHT) = core22;
(*LEFT) = core02;
(*ringOrder) = 9;
break;
case core22:
(*RIGHT) = core32;
(*LEFT) = core12;
(*ringOrder) = 10;
break;
case core32:
(*RIGHT) = core33;
(*LEFT) = core22;
(*ringOrder) = 11;
break;
case core33:
(*RIGHT) = core23;
(*LEFT) = core32;
(*ringOrder) = 12;
break;
case core23:
(*RIGHT) = core13;
(*LEFT) = core33;
(*ringOrder) = 13;
break;
case core13:
(*RIGHT) = core03;
(*LEFT) = core23;
(*ringOrder) = 14;
break;
case core03:
(*RIGHT) = 0x0;
(*LEFT) = core13;
(*ringOrder) = 15;
break;
default: ; // don't know what to do here
}
}