-
Notifications
You must be signed in to change notification settings - Fork 12
/
Copy pathCNTK.py
154 lines (132 loc) · 4.84 KB
/
CNTK.py
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
153
154
import cupy as cp
import numpy as np
import argparse
import scipy.linalg
from utils import load_cifar
parser = argparse.ArgumentParser(description = 'Convolutional Neural Tangent Kernel (CNTK) for CIFAR-10')
parser.add_argument('--depth', default = 21, type = int, help = 'depth of CNTK (#conv layers + 1)')
parser.add_argument('--gap', default = "yes", type = str, help = 'whether GAP (global average pooling) is used')
parser.add_argument('--fix', default = "yes", type = str, help = 'whether first layer and last layer are fixed (or trained) (see Section 4.2 in our paper)')
args = parser.parse_args()
d = args.depth
gap = (args.gap == "yes")
fix = (args.fix == "yes")
#CUDA kernel for convolution operation
conv3 = cp.RawKernel(r'''
extern "C" __global__
void conv3(const float s[32][32][32][32], float t[32][32][32][32])
{
int x1 = threadIdx.x + blockIdx.x - 31;
int y1 = threadIdx.y + blockIdx.y - 31;
int x2 = threadIdx.x;
int y2 = threadIdx.y;
__shared__ float d[32 + 2][32 + 2];
if (x2 == 0){
d[0][y2 + 1] = d[33][y2 + 1] = 0;
if (x2 == 0 && y2 == 0)
d[0][0] = d[0][33] = d[33][0] = d[33][33] = 0;
}
if (y2 == 0){
d[x2 + 1][0] = d[x2 + 1][33] = 0;
}
if (x1 < 0 || x1 > 31 || y1 < 0 || y1 > 31){
d[x2 + 1][y2 + 1] = 0;
return;
}
else
d[x2 + 1][y2 + 1] = s[x1][y1][x2][y2];
__syncthreads();
t[x1][y1][x2][y2] = d[x2][y2] + d[x2][y2 + 1] + d[x2][y2 + 2]
+ d[x2 + 1][y2] + d[x2 + 1][y2 + 1] + d[x2 + 1][y2 + 2]
+ d[x2 + 2][y2] + d[x2 + 2][y2 + 1] + d[x2 + 2][y2 + 2];
}''', 'conv3')
conv_blocks = (63, 63)
conv_threads = (32, 32)
#CUDA kernel for activation
trans = cp.RawKernel(r'''
extern "C" __global__
void trans(float s[32][32][32][32], float t[32][32][32][32], const float l[32][32], const float r[32][32], const float il[32][32], const float ir[32][32])
{
int x1 = blockIdx.x;
int y1 = blockIdx.y;
int x2 = threadIdx.x + ((blockIdx.z >> 2) << 3);
int y2 = threadIdx.y + ((blockIdx.z & 3) << 3);
float S = s[x1][y1][x2][y2], T = t[x1][y1][x2][y2], L = l[x1][y1], R = r[x2][y2], iL = il[x1][y1], iR = ir[x2][y2];
S = S * iL * iR;
float BS = (S * (3.141592654f - acosf(max(min(S, 1.0f), -1.0f))) + sqrtf(1.0f - min(S * S, 1.0f))) * L * R / 28.274333882308138f;
S = (3.141592654f - acosf(max(min(S, 1.0f), -1.0f))) / 28.274333882308138;
t[x1][y1][x2][y2] = T * S + BS;
s[x1][y1][x2][y2] = BS;
}''', 'trans')
trans_blocks = (32, 32, 16)
trans_threads = (8, 8)
#Calculate diagonal entries of $\Sigma^{(h)}(x, x)$ and their reciprocals. See Section 4.3 in our paper.
def xx(x):
RL = [1.0, ]
iRL = [1.0, ]
S = cp.matmul(x.T, x).reshape(32, 32, 32, 32)
conv3(conv_blocks, conv_threads, (S, S))
T = cp.zeros((32, 32, 32, 32), dtype = cp.float32)
if not fix:
T += S
for i in range(1, d - 1):
L = cp.sqrt(cp.diag(S.reshape(1024, 1024)).reshape(32, 32))
iL = 1.0 / L
RL.append(L)
iRL.append(iL)
trans(trans_blocks, trans_threads, (S, T, L, L, iL, iL))
conv3(conv_blocks, conv_threads, (S, S))
conv3(conv_blocks, conv_threads, (T, T))
L = cp.sqrt(cp.diag(S.reshape(1024, 1024)).reshape(32, 32))
iL = 1.0 / L
RL.append(L)
iRL.append(iL)
trans(trans_blocks, trans_threads, (S, T, L, L, iL, iL))
if fix:
T -= S
return RL, iRL
#Caclulate the kernel value of x and z.
#Lx and Lz are diagonal entries of $\Sigma^{(h)}(x, x)$ and $\Sigma^{(h)}(z, z)$.
#iLx and iLz are reciprocals of diagonal entries of $\Sigma^{(h)}(x, x)$ and $\Sigma^{(h)}(z, z)$.
def xz(x, z, Lx, Lz, iLx, iLz):
S = cp.matmul(x.T, z).reshape(32, 32, 32, 32)
conv3(conv_blocks, conv_threads, (S, S))
T = cp.zeros((32, 32, 32, 32), dtype = cp.float32)
if not fix:
T += S
for i in range(1, d - 1):
trans(trans_blocks, trans_threads, (S, T, Lx[i], Lz[i], iLx[i], iLz[i]))
conv3(conv_blocks, conv_threads, (S, S))
conv3(conv_blocks, conv_threads, (T, T))
trans(trans_blocks, trans_threads, (S, T, Lx[-1], Lz[-1], iLx[-1], iLz[-1]))
if fix:
T -= S
return cp.mean(T) if gap else cp.trace(T.reshape(1024, 1024))
#Load CIFAR-10.
(X_train, y_train), (X_test, y_test) = load_cifar()
X = np.concatenate((X_train, X_test), axis = 0)
N = X.shape[0]
N_train = X_train.shape[0]
N_test = X_test.shape[0]
X = cp.asarray(X).reshape(-1, 3, 1024)
#Calculate diagonal entries.
L = []
iL = []
for i in range(N):
Lx, iLx = xx(X[i])
L.append(Lx)
iL.append(iLx)
#####Calculate kernel values.
#####Below we provide a naive implementation using for-loops.
#####Parallelize this part according to your specific computing enviroment to utilize multiple GPUs.
H = np.zeros((N, N), dtype = np.float32)
for i in range(N):
for j in range(N):
H[i][j] = xz(X[i], X[j], L[i], L[j], iL[i], iL[j])
#####
#Solve kernel regression.
Y_train = np.ones((N_train, 10)) * -0.1
for i in range(N_train):
Y_train[i][y_train[i]] = 0.9
u = H[N_train:, :N_train].dot(scipy.linalg.solve(H[:N_train, :N_train], Y_train))
print "test accuracy:", 1.0 * np.sum(np.argmax(u, axis = 1) == y_test) / N_test