 	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) {

	}
}

__global__ void preact_c1(float input[28][28], float preact[6][24][24], float weight[6][5][5])
{
	const int pos = blockIdx.x * blockDim.x + threadIdx.x;
	const int size = blockDim.x * gridDim.x;

	const int N = 5*5*6*24*24;

	for (int n = N * pos / size; n < N * (pos+1) / size; ++n) {
		int idx = n;
		const int i1 = ((idx /= 1	) % 5);
		const int i2 = ((idx /= 5	) % 5);
		const int i3 = ((idx /= 5	) % 6);
		const int i4 = ((idx /= 6	) % 24);
		const int i5 = ((idx /= 24	) % 24);

		atomicAdd(&preact[i3][i4][i5], weight[i3][i1][i2] * input[i4 + i1][i5 + i2]);
	}
}

__global__ void bias_c1(float preact[6][24][24], float bias[6])
{
	const int pos = blockIdx.x * blockDim.x + threadIdx.x;
	const int size = blockDim.x * gridDim.x;

	const int N = 6*24*24;

	for (int n = N * pos / size; n < N * (pos+1) / size; ++n) {
		int idx = n;
		const int i1 = ((idx /= 1	) % 6);
		const int i2 = ((idx /= 6	) % 24);
		const int i3 = ((idx /= 24	) % 24);

		preact[i1][i2][i3] += bias[i1];
	}
}

__global__ void preact_s1(float input[6][24][24], float preact[6][6][6], float weight[1][4][4])
{
	const int pos = blockIdx.x * blockDim.x + threadIdx.x;
	const int size = blockDim.x * gridDim.x;

	const int N = 4*4*6*6*6;

	for (int n = N * pos / size; n < N * (pos+1) / size; ++n) {
		int idx = n;
		const int i1 = ((idx /= 1	) % 4);
		const int i2 = ((idx /= 4	) % 4);
		const int i3 = ((idx /= 4	) % 6);
		const int i4 = ((idx /= 6	) % 6);
		const int i5 = ((idx /= 6	) % 6);

		atomicAdd(&preact[i3][i4][i5], weight[0][i1][i2] * input[i3][i4 * 4 + i1][i5 * 4 + i2]);
	}
}

__global__ void bias_s1(float preact[6][6][6], float bias[1])
{
	const int pos = blockIdx.x * blockDim.x + threadIdx.x;
	const int size = blockDim.x * gridDim.x;

	const int N = 6*6*6;

	for (int n = N * pos / size; n < N * (pos+1) / size; ++n) {
		int idx = n;
		const int i1 = ((idx /= 1	) % 6);
		const int i2 = ((idx /= 6	) % 6);
		const int i3 = ((idx /= 6	) % 6);

		preact[i1][i2][i3] += bias[0];
	}
}

__global__ void preact_f(float input[6][6][6], float preact[10], float weight[10][6][6][6])
{
	const int pos = blockIdx.x * blockDim.x + threadIdx.x;
	const int size = blockDim.x * gridDim.x;

	const int N = 10*6*6*6;

	for (int n = N * pos / size; n < N * (pos+1) / size; ++n) {
		int idx = n;
		const int i1 = ((idx /= 1	) % 10);
		const int i2 = ((idx /= 10	) % 6);
		const int i3 = ((idx /= 6	) % 6);
		const int i4 = ((idx /= 6	) % 6);

		atomicAdd(&preact[i1], weight[i1][i2][i3][i4] * input[i2][i3][i4]);
	}
}

__global__ void bias_f(float preact[10], float bias[10])
{
	const int pos = blockIdx.x * blockDim.x + threadIdx.x;
	const int size = blockDim.x * gridDim.x;

	const int N = 10;

	for (int idx = N * pos / size; idx < N * (pos+1) / size; ++idx) {

			l_f.bp_clear();
			l_s1.bp_clear();
			l_c1.bp_clear();
			l_s2.bp_clear();
			l_c2.bp_clear();

			makeError<<<10, 1>>>(l_f.d_preact, l_f.output, train_set[i].label, 10);

			bp_weight_f<<<64, 64>>>((float (*)[36][4][4])l_f.d_weight, l_f.d_preact, (float (*)[4][4])l_s2.output);
			bp_bias_f<<<64, 64>>>(l_f.bias, l_f.d_preact);

			bp_output_s2<<<64, 64>>>((float (*)[4][4])l_s2.d_output, (float (*)[36][4][4])l_f.weight, l_f.d_preact);
			BP_PREACT_S(36, 4, l_s2.d_preact, l_s2.d_output, l_s2.preact);
			BP_WEIGHT_S(2, 36, 4, l_s2.d_weight, l_s2.d_preact, l_c2.output);
			BP_BIAS_S(36, 4, l_s2.bias, l_s2.d_preact);

			BP_OUTPUT_C(2, 36, 4, l_c2.d_output, l_s2.weight, l_s2.d_preact);
			BP_PREACT_C(36, 8, l_c2.d_preact, l_c2.d_output, l_c2.preact);
			BP_WEIGHT_C(6, 5, 8, 6, l_c2.d_weight, l_c2.d_preact, l_s2.output);
			BP_BIAS_C(6, 8, 6, l_c2.bias, l_c2.preact);

			BP_OUTPUT_S(6, 12, 5, 6, l_s1.d_output, l_c2.weight, l_c2.d_preact);
			BP_PREACT_S(6, 12, l_s1.d_preact, l_s1.d_output, l_s1.preact);
			BP_WEIGHT_S(2, 6, 12, l_s1.d_weight, l_s1.d_preact, l_c1.output);
			BP_BIAS_S(6, 12, l_s1.bias, l_s1.d_preact);

			BP_OUTPUT_C(2, 6, 12, l_c1.d_output, l_s1.weight, l_s1.d_preact);
			BP_PREACT_C(6, 24, l_c1.d_preact, l_c1.d_output, l_c1.preact);
			BP_WEIGHT_C(6, 5, 23, 1, l_c1.d_weight, l_c1.d_preact, l_input.output);
			BP_BIAS_C(6, 24, 1, l_c1.bias, l_c1.preact);

			apply_grad<<<64, 64>>>(l_f.weight, l_f.d_weight, l_f.M * l_f.N);
			apply_grad<<<64, 64>>>(l_s2.weight, l_s2.d_weight, l_s2.M * l_s2.N);
			apply_grad<<<64, 64>>>(l_c2.weight, l_c2.d_weight, l_c2.M * l_c2.N);
			apply_grad<<<64, 64>>>(l_s1.weight, l_s1.d_weight, l_s1.M * l_s1.N);
			apply_grad<<<64, 64>>>(l_c1.weight, l_c1.d_weight, l_c1.M * l_c1.N);
		}
