-
Notifications
You must be signed in to change notification settings - Fork 0
/
analysis.cu
153 lines (111 loc) · 3.75 KB
/
analysis.cu
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
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
#include <stdlib.h>
#include <time.h>
#define CA_SIZE %s
#define CA_ITERATIONS %s
#define CA_REPEAT %s
#define CONNECTION_RADIUS %s
#define N_CONNECTIONS %s
#define N_OBSERVATIONS %s
#define K_HISTORY %s
#define N_POSSIBLE_HISTORY %s
__global__ void kernel_calc_diffs(unsigned int majority[CA_REPEAT],
unsigned int executions[CA_REPEAT][CA_ITERATIONS + 1][CA_SIZE],
unsigned int sum_diffs[CA_SIZE]) {
int cell = threadIdx.x;
int errors_local[N_CONNECTIONS];
int sum_diffs_local;
int i, repeat, shift;
__shared__ int errors[N_CONNECTIONS];
__shared__ int flow;
for (i = 0; i < N_CONNECTIONS; i++)
errors_local[i] = 0;
if (cell == 0)
for (i = 0; i < N_CONNECTIONS; i++)
errors[i] = 0;
__syncthreads();
for (repeat = 0; repeat < CA_REPEAT; repeat++)
for (i = 0; i < CA_ITERATIONS; i++)
for (shift = -CONNECTION_RADIUS; shift <= CONNECTION_RADIUS;
shift++)
errors_local[shift + CONNECTION_RADIUS] +=
abs((float) executions[repeat][i][cell] -
executions[repeat][i + 1][(cell + shift + CA_SIZE) %%
CA_SIZE]);
for (i = 0; i < N_CONNECTIONS; i++)
atomicAdd(&errors[i], errors_local[i]);
__syncthreads();
if (cell == 0) {
int min_error = errors[0];
flow = -CONNECTION_RADIUS;
for (i = 1; i < N_CONNECTIONS; i++)
if (errors[i] < min_error) {
min_error = errors[i];
flow = i - CONNECTION_RADIUS;
}
}
__syncthreads();
for (repeat = 0; repeat < CA_REPEAT; repeat++) {
sum_diffs_local = 0;
for (i = 0; i < CA_ITERATIONS; i++) {
int diff = abs((float) executions[repeat][i + 1][cell] -
executions[repeat][i][(cell - flow + CA_SIZE) %% CA_SIZE]);
sum_diffs_local += diff;
}
sum_diffs[cell] += sum_diffs_local;
}
}
__device__ int to_binary(unsigned int execution[CA_ITERATIONS + 1][CA_SIZE],
int cell, int i) {
int k;
int value = 0;
for (k = 0; k < K_HISTORY; k++)
value += execution[i - k][cell] << k;
return value;
}
__global__ void kernel_probabilities(
unsigned int execution[CA_REPEAT][CA_ITERATIONS + 1][CA_SIZE],
float p_joint_table[CA_SIZE][N_POSSIBLE_HISTORY][2],
float p_prev_table[CA_SIZE][N_POSSIBLE_HISTORY],
float p_curr_table[CA_SIZE][2]) {
int cell = threadIdx.x;
int repeat = blockIdx.x;
int i;
for (i = K_HISTORY; i < CA_ITERATIONS; i++) {
int past = to_binary(execution[repeat], cell, i - 1);
int present = execution[repeat][i][cell];
atomicAdd(&p_joint_table[cell][past][present], 1);
atomicAdd(&p_prev_table[cell][past], 1);
atomicAdd(&p_curr_table[cell][present], 1);
}
}
__global__ void kernel_active_storage(
unsigned int execution[CA_REPEAT][K_HISTORY + 1][CA_SIZE],
float p_joint_table[CA_SIZE][N_POSSIBLE_HISTORY][2],
float p_prev_table[CA_SIZE][N_POSSIBLE_HISTORY],
float p_curr_table[CA_SIZE][2], float active_storage[CA_SIZE]) {
int cell = threadIdx.x;
int repeat = blockIdx.x;
int past, present;
float p1, p2, p3;
past = to_binary(execution[repeat], cell, K_HISTORY - 1);
present = execution[repeat][K_HISTORY][cell];
p1 = p_joint_table[cell][past][present] / (float) N_OBSERVATIONS;
p2 = p_prev_table[cell][past] / (float) N_OBSERVATIONS;
p3 = p_curr_table[cell][present] / (float) N_OBSERVATIONS;
atomicAdd(&active_storage[cell], log2f(p1 / (p2 * p3)) / N_OBSERVATIONS);
}
__global__ void kernel_entropy_rate(
float p_joint_table[CA_SIZE][N_POSSIBLE_HISTORY][2],
float p_prev_table[CA_SIZE][N_POSSIBLE_HISTORY],
float entropy_rate[CA_SIZE]) {
int cell = threadIdx.x;
int i, j;
float aux = 0;
for (i = 0; i < N_POSSIBLE_HISTORY; i++)
for (j = 0; j < 2; j++)
if (p_joint_table[cell][i][j] > 0)
aux -= p_joint_table[cell][i][j] *
log2f((float) p_joint_table[cell][i][j] /
p_prev_table[cell][i]) / N_OBSERVATIONS;
entropy_rate[cell] = aux;
}