 			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 (*)[6][6][6])l_f.d_weight, l_f.d_preact, (float (*)[6][6])l_s1.output);
+			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_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);
-
+			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);
 		}

	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);
}

	cublasCreate(&blas);

	float err;
	int iter = 3;

	while (iter < 0 || iter-- > 0) {
		err = 0.0f;
		
		for (int i = 0; i < train_cnt; ++i) {
