 	cublasCreate(&blas);
 
 	float err;
-
-	while (true) {
+	int iter = 3;
+
+	while (iter < 0 || iter-- > 0) {
 		err = 0.0f;
 		
 		for (int i = 0; i < train_cnt; ++i) {

	l_input.clear();
	l_c1.clear();
	l_s1.clear();
	l_f.clear();

	l_input.setOutput((float *)input);
	
	preact_c1<<<64, 64>>>((float (*)[28])l_input.output, (float (*)[24][24])l_c1.preact, (float (*)[5][5])l_c1.weight);
	bias_c1<<<64, 64>>>((float (*)[24][24])l_c1.preact, l_c1.bias);
	apply_step_function<<<64, 64>>>(l_c1.preact, l_c1.output, l_c1.O);

	preact_s1<<<64, 64>>>((float (*)[24][24])l_c1.output, (float (*)[6][6])l_s1.preact, (float (*)[4][4])l_s1.weight);
	bias_s1<<<64, 64>>>((float (*)[6][6])l_s1.preact, l_s1.bias);
	apply_step_function<<<64, 64>>>(l_s1.preact, l_s1.output, l_s1.O);

	preact_f<<<64, 64>>>((float (*)[6][6])l_s1.output, l_f.preact, (float (*)[6][6][6])l_f.weight);
	bias_f<<<64, 64>>>(l_f.preact, l_f.bias);
	apply_step_function<<<64, 64>>>(l_f.preact, l_f.output, l_f.O);
}

};

static cublasHandle_t blas;
static Layer l_input(0, 0, 28*28), l_c1(5*5, 6, 6*24*24), l_s1(2*2, 1, 6*12*12),
	     l_c2(5*5, 6, 36*8*8), l_s2(2*2, 1, 36*4*4), l_f(36*4*4, 10, 10);

#define FOR6(I1, I2, I3, I4, I5, I6, M) \
	const int pos = blockIdx.x * blockDim.x + threadIdx.x; \
	const int size = blockDim.x * gridDim.x; \
	\
	const int N = I1*I2*I3*I4*I5*I6; \
	\
	for (int n = N * pos / size; n < N * (pos+1) / size; ++n) { \
		int idx = n; \
		const int i1 = ((idx /= 1	) % I1); \
		const int i2 = ((idx /= I1	) % I2); \
		const int i3 = ((idx /= I2	) % I3); \
		const int i4 = ((idx /= I3	) % I4); \
		const int i5 = ((idx /= I4	) % I5); \
		const int i6 = ((idx /= I5	) % I6); \
		\
		M; \
	} \

#define FOR5(I1, I2, I3, I4, I5, M)	FOR6(I1, I2, I3, I4, I5, 1, M)
#define FOR4(I1, I2, I3, I4, M)		FOR6(I1, I2, I3, I4, 1, 1, M)
#define FOR3(I1, I2, I3, M)		FOR6(I1, I2, I3, 1, 1, 1, M)
#define FOR2(I1, I2, M)			FOR6(I1, I2, 1, 1, 1, 1, M)
#define FOR1(I1, M)			FOR6(I1, 1, 1, 1, 1, 1, M)

#define CAST0(I, v)	((float (*)[I])v)
#define CAST1(I, v)	((float (*)[I][I])v)
#define CAST2(V, I, v)	((float (*)[V][I][I])v)

#define PREACT_C(I1, I2, I3, I4, input, preact, weight) \
	preact_c<I1, I2, I3, I4><<<64, 64>>>((float (*)[I3+I1-1][I3+I1-1])input, (float (*)[I3][I3])preact, (float (*)[I1][I1])weight)

#define BIAS_C(I1, I2, preact, bias) \
	bias_c<I1, I2><<<64, 64>>>((float (*)[I2][I2])preact, bias);

#define PREACT_S(I1, I2, I3, input, preact, weight) \
	preact_s<I1, I2, I3><<<64, 64>>>((float (*)[I1*I3][I1*I3])input, (float (*)[I3][I3])preact, (float (*)[I1])weight)

#define BIAS_S(I1, I2, preact, bias) \
	bias_s<I1, I2><<<64, 64>>>((float (*)[I2][I2])preact, bias)

#define BP_OUTPUT_S(I1, I2, I3, I4, d_output, n_weight, nd_preact) \
	bp_output_s<I1, I2, I3, I4><<<64, 64>>>(CAST1(I2+I3-1, d_output), CAST1(I2, n_weight), CAST1(I3, nd_preact))

#define BP_PREACT_S(I1, I2, d_preact, d_output, preact) \
	bp_preact_s<I1, I2><<<64, 64>>>((float (*)[I2][I2])d_preact, (float (*)[I2][I2])d_output, (float (*)[I2][I2])preact)

#define BP_WEIGHT_S(I1, I2, I3, d_weight, d_preact, p_output) \
	bp_weight_s<I1, I2, I3><<<64, 64>>>(CAST0(I1, d_preact), CAST1(I3, d_preact), CAST1(I1*I3, p_output))

#define BP_BIAS_S(I1, I2, bias, d_preact) \
	bp_bias_s<I1, I2><<<64, 64>>>(bias, CAST1(I2, d_preact))


#define BP_OUTPUT_C(I1, I2, I3, d_output, n_weight, nd_preact) \
	bp_output_c<I1, I2, I3><<<64, 64>>>(CAST1(I1*I3, d_output), CAST0(I1, n_weight), CAST1(I3, nd_preact))

#define BP_PREACT_C(I1, I2, d_preact, d_output, preact) \
	bp_preact_c<I1, I2><<<64, 64>>>(CAST1(I2, d_preact), CAST1(I2, d_output), CAST1(I2, preact))

#define BP_WEIGHT_C(I1, I2, I3, I4, d_weight, d_preact, p_output) \
	bp_weight_c<I1, I2, I3, I4><<<64, 64>>>(CAST1(I2, d_weight), CAST1(I3, d_preact), CAST1(I3 + I2 - 1, p_output))

#define BP_BIAS_C(I1, I2, I3, bias, d_preact) \
	bp_bias_c<I1, I2, I3><<<64, 64>>>(bias, CAST1(I2, d_preact))

__device__ float step_function(float v)
{
