 			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_weight_f<<<64, 64>>>((float (*)[6][6][6])l_f.d_weight, l_f.d_preact, (float (*)[6][6])l_s1.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);
+			bp_output_s1<<<64, 64>>>((float (*)[6][6])l_s1.d_output, (float (*)[6][6][6])l_f.weight, l_f.d_preact);
+			bp_preact_s1<<<64, 64>>>((float (*)[6][6])l_s1.d_preact, (float (*)[6][6])l_s1.d_output, (float (*)[6][6])l_s1.preact);
+			bp_weight_s1<<<64, 64>>>((float (*)[4][4])l_s1.d_weight, (float (*)[6][6])l_s1.d_preact, (float (*)[24][24])l_c1.output);
+			bp_bias_s1<<<64, 64>>>(l_s1.bias, (float (*)[6][6])l_s1.d_preact);
+
+			bp_output_c1<<<64, 64>>>((float (*)[24][24])l_c1.d_output, (float (*)[4][4])l_s1.weight, (float (*)[6][6])l_s1.d_preact);
+			bp_preact_c1<<<64, 64>>>((float (*)[24][24])l_c1.d_preact, (float (*)[24][24])l_c1.d_output, (float (*)[24][24])l_c1.preact);
+			bp_weight_c1<<<64, 64>>>((float (*)[5][5])l_c1.d_weight, (float (*)[24][24])l_c1.d_preact, (float (*)[28])l_input.output);
+			bp_bias_c1<<<64, 64>>>(l_c1.bias, (float (*)[24][24])l_c1.d_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);
 		}

	}
}

__global__ void bp_weight_f(float d_weight[10][36][4][4], float d_preact[10], float p_output[36][4][4])
{
	FOR4(10, 36, 4, 4, d_weight[i1][i2][i3][i4] = d_preact[i1] * p_output[i2][i3][i4]);
}

__global__ void bp_bias_f(float bias[10], float d_preact[10])
{
	FOR1(10, bias[i1] += dt * d_preact[i1]);
}

__global__ void bp_output_s2(float d_output[36][4][4], float n_weight[10][36][4][4], float nd_preact[10])
{
	FOR4(10, 36, 4, 4, atomicAdd(&d_output[i2][i3][i4], n_weight[i1][i2][i3][i4] * nd_preact[i1]));
}

template<int I1, int I2, int I3, int I4>
__global__ void bp_output_s(float d_output[I1][I2+I3-1][I2+I3-1], float n_weight[I4][I2][I2], float nd_preact[I1*I4][I3][I3])
{
	FOR6(I1, I2, I2, I3, I3, I4,
		atomicAdd(&d_output[i1][i2+i4-1][i3+i5-1], n_weight[i6][i2][i3] * nd_preact[i6 * I1 + i1][i4][i5])
	);
}

template<int I1, int I2>
__global__ void bp_preact_s(float d_preact[I1][I2][I2], float d_output[I1][I2][I2], float preact[I1][I2][I2])
{
	FOR3(I1, I2, I2, 
		const float o = step_function(preact[i1][i2][i3]);

		d_preact[i1][i2][i3] = d_output[i1][i2][i3] * o * (1 - o);
	);
}

template<int I1, int I2, int I3>
__global__ void bp_weight_s(float d_weight[I1][I1], float d_preact[I2][I3][I3], float p_output[I2][I1*I3][I1*I3])
{
	const float d = I2 * I3 * I3;

	FOR5(I1, I1, I2, I3, I3,
		atomicAdd(&d_weight[i1][i2], d_preact[i3][i4][i5] * p_output[i3][i4 * I1 + i1][i5 * I1 + i2]) / d);
}

template<int I1, int I2>
__global__ void bp_bias_s(float bias[1], float d_preact[I1][I2][I2])
{
	const float d = I1 * I2 * I2;

	FOR3(I1, I2, I2, atomicAdd(&bias[0], dt * d_preact[i1][i2][i3] / d));
}

template<int I1, int I2, int I3>
__global__ void bp_output_c(float d_output[I2][I1*I3][I1*I3], float n_weight[I1][I1], float nd_preact[I2][I3][I3])
{
	FOR5(I1, I1, I2, I3, I3,
		atomicAdd(&d_output[i3][i4 * I1 + i1][i5 * I1 + i2], n_weight[i1][i2] * nd_preact[i3][i4][i5]));
}

template<int I1, int I2>
__global__ void bp_preact_c(float d_preact[I1][I2][I2], float d_output[I1][I2][I2], float preact[I1][I2][I2])
{
	FOR3(I1, I2, I2, 
		const float o = step_function(preact[i1][i2][i3]);

		d_preact[i1][i2][i3] = d_output[i1][i2][i3] * o * (1 - o);
	);
}

template<int I1, int I2, int I3, int I4>
__global__ void bp_weight_c(float d_weight[I1][I2][I2], float d_preact[I1*I4][I3][I3],
	float p_output[I4][I3 + I2 - 1][I3 + I2 - 1])
{
	const float d = I3 * I3 * I4;

	FOR6(I1, I2, I2, I3, I3, I4,
		atomicAdd(&d_weight[i1][i2][i3], d_preact[i6 * I1 + i1][i4][i5] * p_output[i6][i4 + i2][i5 + i3] / d));
}

template<int I1, int I2, int I3>
__global__ void bp_bias_c(float bias[I1], float d_preact[I1 * I3][I2][I2])
{
	const float d = I2 * I2 * I3;

	FOR4(I1, I2, I2, I3, atomicAdd(&bias[i1], dt * d_preact[i1 * I3 + i4][i2][i3] / d));
}

static void propagate(double data[28][28])

#include <cublas_v2.h>

const static float dt = 1.0E-01f;
const static float threshold = 1.0E-02f;

struct Layer {
	float *output;
